From d23c2e8cd2e1f242d74cda50dc0867cd070856b9 Mon Sep 17 00:00:00 2001 From: Xin Jin Date: Mon, 12 May 2025 14:58:26 +0100 Subject: [PATCH 1/4] Fix samples memory leak and cleanup issues Fix memory leaks and cleanup logic in the following sample programs: - samples/core/binaries - samples/core/blur - samples/core/reduce - samples/core/saxpy - samples/core/multi-device - samples/extensions/khr/externalmemory Also fix minor logic and error-handling inconsistencies such as: - Ensuring proper release of cl_event objects - Correct ordering of Vulkan cleanup calls - Improved version string checks and CLI option allocation Signed-off-by: Xin Jin --- samples/core/binaries/main.c | 8 +- samples/core/blur/main.c | 93 ++++++++++++------- samples/core/multi-device/main.c | 21 +++-- samples/core/reduce/main.c | 11 ++- samples/core/saxpy/main.c | 14 +-- samples/extensions/khr/externalmemory/main.c | 20 ++-- .../extensions/khr/externalmemory/main.cpp | 10 +- 7 files changed, 111 insertions(+), 66 deletions(-) diff --git a/samples/core/binaries/main.c b/samples/core/binaries/main.c index c1a09abc..15da57ae 100644 --- a/samples/core/binaries/main.c +++ b/samples/core/binaries/main.c @@ -258,21 +258,21 @@ int main(int argc, char *argv[]) error, vec); /// Run kernel - cl_event pass; OCLERROR_RET(clSetKernelArg(Collatz, 0, sizeof(cl_mem), &buf), error, buff); GET_CURRENT_TIMER(start_time) + cl_event pass; OCLERROR_RET(clEnqueueNDRangeKernel(queue, Collatz, 1, &start, &length, NULL, 0, NULL, &pass), error, buff); - OCLERROR_RET(clWaitForEvents(1, &pass), error, buff); + OCLERROR_RET(clWaitForEvents(1, &pass), error, ev); GET_CURRENT_TIMER(end_time) if (diag_opts.verbose) print_timings(start_time, end_time, &pass, 1); OCLERROR_RET(clEnqueueReadBuffer(queue, buf, CL_BLOCKING, 0, sizeof(cl_int) * length, v, 0, NULL, NULL), - error, buff); + error, ev); /// Show results int max_steps = 0; @@ -298,6 +298,8 @@ int main(int argc, char *argv[]) length, start + 1, max_steps, max_ind); /// Cleanup +ev: + OCLERROR_RET(clReleaseEvent(pass), end_error, buff); buff: OCLERROR_RET(clReleaseMemObject(buf), end_error, vec); vec: diff --git a/samples/core/blur/main.c b/samples/core/blur/main.c index 2d1a4d38..025af91e 100644 --- a/samples/core/blur/main.c +++ b/samples/core/blur/main.c @@ -344,7 +344,7 @@ cl_int single_pass_box_blur(state *const s, cl_int size) // compile kernel cl_kernel blur; OCLERROR_PAR(blur = clCreateKernel(s->program, "blur_box", &error), error, - end); + endl); // set kernel parameters OCLERROR_RET(clSetKernelArg(blur, 0, sizeof(cl_mem), &s->input_image_buf), @@ -359,22 +359,23 @@ cl_int single_pass_box_blur(state *const s, cl_int size) OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur, 2, origin, image_size, NULL, 0, NULL, &pass), error, blr); - OCLERROR_RET(clWaitForEvents(1, &pass), error, blr); + OCLERROR_RET(clWaitForEvents(1, &pass), error, ev); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr); + error, ev); if (s->verbose) print_timings(start, end, &pass, 1); // write output file - OCLERROR_RET(finalize_blur(s), error, blr); - + OCLERROR_RET(finalize_blur(s), error, ev); +ev: + clReleaseEvent(pass); blr: clReleaseKernel(blur); -end: +endl: return error; } @@ -417,20 +418,23 @@ cl_int dual_pass_box_blur(state *const s, cl_int size) error, blr2); OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, image_size, NULL, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); - + OCLERROR_RET(finalize_blur(s), error, ev2); +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -531,20 +535,23 @@ cl_int dual_pass_local_memory_exchange_box_blur(state *const s, cl_int size) size_t wgss[3] = { 1, wgs2, 1 }; OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, work_size2, wgss, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); - + OCLERROR_RET(finalize_blur(s), error, ev2); +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -620,21 +627,25 @@ cl_int dual_pass_subgroup_exchange_box_blur(state *const s, cl_int size) size_t wgss[3] = { 1, wgs2, 1 }; OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, work_size2, wgss, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); + OCLERROR_RET(finalize_blur(s), error, ev2); // cleanup for error handling +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -685,20 +696,23 @@ cl_int dual_pass_kernel_blur(state *const s, cl_int size, cl_mem kern) error, blr2); OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, image_size, NULL, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); - + OCLERROR_RET(finalize_blur(s), error, ev2); +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -801,20 +815,23 @@ cl_int dual_pass_local_memory_exchange_kernel_blur(state *const s, cl_int size, size_t wgss[3] = { 1, wgs2, 1 }; OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, work_size2, wgss, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); - + OCLERROR_RET(finalize_blur(s), error, ev2); +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -894,21 +911,25 @@ cl_int dual_pass_subgroup_exchange_kernel_blur(state *const s, cl_int size, size_t wgss[3] = { 1, wgs2, 1 }; OCLERROR_RET(clEnqueueNDRangeKernel(s->queue, blur2, 2, origin, work_size2, wgss, 0, NULL, pass + 1), - error, blr2); - OCLERROR_RET(clWaitForEvents(2, pass), error, blr2); + error, ev1); + OCLERROR_RET(clWaitForEvents(2, pass), error, ev2); GET_CURRENT_TIMER(end) OCLERROR_RET(clEnqueueReadImage(s->queue, s->output_image_buf, CL_BLOCKING, origin, image_size, 0, 0, s->output_image.pixels, 0, NULL, NULL), - error, blr2); + error, ev2); if (s->verbose) print_timings(start, end, pass, 2); // write output file - OCLERROR_RET(finalize_blur(s), error, blr2); + OCLERROR_RET(finalize_blur(s), error, ev2); // cleanup for error handling +ev2: + OCLERROR_RET(clReleaseEvent(pass[1]), end_error, ev1); +ev1: + OCLERROR_RET(clReleaseEvent(pass[0]), end_error, blr2); blr2: OCLERROR_RET(clReleaseKernel(blur2), end_error, blr1); blr1: @@ -1073,11 +1094,11 @@ int main(int argc, char *argv[]) dev_version_size, dev_version, NULL), error, dev); char compiler_options[1024] = ""; - if (opencl_version_contains(dev_version, "3.")) + if (opencl_version_contains(dev_version, "OpenCL C 3.")) { strcat(compiler_options, "-cl-std=CL3.0 "); } - else if (opencl_version_contains(dev_version, "2.")) + else if (opencl_version_contains(dev_version, "OpenCL C 2.")) { strcat(compiler_options, "-cl-std=CL2.0 "); } diff --git a/samples/core/multi-device/main.c b/samples/core/multi-device/main.c index 57cc573e..2bf2a856 100644 --- a/samples/core/multi-device/main.c +++ b/samples/core/multi-device/main.c @@ -208,12 +208,6 @@ int main(int argc, char* argv[]) dev_opts.triplet.dev_type, &error), error, end); - // Query OpenCL version supported by device. - char dev_version[64]; - OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_VERSION, sizeof(dev_version), - &dev_version, NULL), - error, end); - if (!diag_opts.quiet) { cl_util_print_device_info(dev); @@ -225,6 +219,19 @@ int main(int argc, char* argv[]) fflush(stdout); } + // Query OpenCL version supported by device. + size_t dev_version_size; + OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_VERSION, 0, NULL, + &dev_version_size), + error, end); + + char *dev_version = malloc(dev_version_size); + if (!dev_version) + { + fprintf(stderr, "Failed to allocate memory for device version.\n"); + exit(EXIT_FAILURE); + } + if (opencl_version_contains(dev_version, "1.0") || opencl_version_contains(dev_version, "1.1")) { @@ -233,8 +240,10 @@ int main(int argc, char* argv[]) "1.2 feature, but the device chosen only supports OpenCL %s. " "Please try with a different OpenCL device instead.\n", dev_version); + free(dev_version); exit(EXIT_SUCCESS); } + free(dev_version); // Check if device supports fission. cl_device_partition_property* dev_props = NULL; diff --git a/samples/core/reduce/main.c b/samples/core/reduce/main.c index b4acbd31..3507941f 100644 --- a/samples/core/reduce/main.c +++ b/samples/core/reduce/main.c @@ -457,7 +457,7 @@ int main(int argc, char *argv[]) size_t gl = global(curr, factor, wgs); OCLERROR_RET(clEnqueueNDRangeKernel(queue, reduce, 1, NULL, &gl, &wgs, 0, NULL, pass), - error, pas); + error, ev); curr = new_size(curr, factor); ++pass; @@ -469,7 +469,7 @@ int main(int argc, char *argv[]) } } - OCLERROR_RET(clWaitForEvents(steps, passes), error, pas); + OCLERROR_RET(clWaitForEvents(steps, passes), error, ev); GET_CURRENT_TIMER(dev_end) cl_ulong dev_time; @@ -488,7 +488,7 @@ int main(int argc, char *argv[]) OCLERROR_RET(clEnqueueReadBuffer(queue, back, CL_BLOCKING, 0, sizeof(cl_int), (void *)&dev_res, 0, NULL, NULL), - error, pas); + error, ev); // Validate if (dev_res != seq_ref) @@ -516,6 +516,11 @@ int main(int argc, char *argv[]) (unsigned long long)(host_time + 500) / 1000); } +ev: + for (cl_event *pass_release = passes; pass_release < pass; ++pass_release) + { + OCLERROR_RET(clReleaseEvent(*pass_release), end_error, pas); + } pas: free(passes); bufb: diff --git a/samples/core/saxpy/main.c b/samples/core/saxpy/main.c index 92c6d0ee..4a2f38ce 100644 --- a/samples/core/saxpy/main.c +++ b/samples/core/saxpy/main.c @@ -62,21 +62,19 @@ cl_int parse_options(int argc, char *argv[], struct options_Saxpy *saxpy_opts) { cl_int error = CL_SUCCESS; - struct cag_option *opts = NULL, *tmp = NULL; + struct cag_option *opts = NULL; size_t n = 0; /* Prepare all options array. */ MEM_CHECK(opts = add_CLI_options(opts, &n, DiagnosticOptions, CAG_ARRAY_SIZE(DiagnosticOptions)), error, end); - MEM_CHECK(tmp = add_CLI_options(opts, &n, SingleDeviceOptions, - CAG_ARRAY_SIZE(SingleDeviceOptions)), + MEM_CHECK(opts = add_CLI_options(opts, &n, SingleDeviceOptions, + CAG_ARRAY_SIZE(SingleDeviceOptions)), error, end); - opts = tmp; - MEM_CHECK(tmp = add_CLI_options(opts, &n, SaxpyOptions, - CAG_ARRAY_SIZE(SaxpyOptions)), + MEM_CHECK(opts = add_CLI_options(opts, &n, SaxpyOptions, + CAG_ARRAY_SIZE(SaxpyOptions)), error, end); - opts = tmp; char identifier; cag_option_context cag_context; @@ -143,8 +141,6 @@ int main(int argc, char *argv[]) OCLERROR_PAR(context = clCreateContext(NULL, 1, &device, NULL, NULL, &error), error, end); - OCLERROR_PAR(queue = clCreateCommandQueue(context, device, 0, &error), - error, cont); OCLERROR_RET(clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &platform, NULL), error, cont); diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 0ab03568..4bd5f201 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -312,9 +312,6 @@ int main(int argc, char* argv[]) cl_context_properties context_props[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)cl_platform, 0 }; - OCLERROR_PAR(context = clCreateContext(context_props, 1, &cl_device, NULL, - NULL, &error), - error, end); // Check if the device supports the Khronos extensions needed before // attempting to compile the kernel. @@ -330,9 +327,15 @@ int main(int argc, char* argv[]) fprintf(stdout, "OpenCL device does not support the required Khronos " "extensions\n"); + vkDestroyDevice(vk_device, NULL); + vkDestroyInstance(instance, NULL); exit(EXIT_SUCCESS); } + OCLERROR_PAR(context = clCreateContext(context_props, 1, &cl_device, NULL, + NULL, &error), + error, vk); + // Compile kernel. if (diag_opts.verbose) { @@ -625,7 +628,7 @@ int main(int argc, char* argv[]) OCLERROR_RET(clEnqueueNDRangeKernel(queue, saxpy, 1, NULL, &length, &wgs, 0, NULL, &kernel_run), error, que); - OCLERROR_RET(clWaitForEvents(1, &kernel_run), error, que); + OCLERROR_RET(clWaitForEvents(1, &kernel_run), error, ev); GET_CURRENT_TIMER(dev_end) cl_ulong dev_time; @@ -661,7 +664,7 @@ int main(int argc, char* argv[]) OCLERROR_RET(clEnqueueReadBuffer(queue, cl_buf_y, CL_BLOCKING, 0, sizeof(cl_float) * length, (void*)arr_x, 0, NULL, NULL), - error, que); + error, ev); // Validate solution. for (size_t i = 0; i < length; ++i) @@ -693,6 +696,8 @@ int main(int argc, char* argv[]) } // Release resources. +ev: + OCLERROR_RET(clReleaseEvent(kernel_run), end_error, que); que: OCLERROR_RET(clReleaseCommandQueue(queue), end_error, cont); clbufy: @@ -717,7 +722,10 @@ int main(int argc, char* argv[]) ker: free(kernel); cont: - OCLERROR_RET(clReleaseContext(context), end_error, end); + OCLERROR_RET(clReleaseContext(context), end_error, vk); +vk: + vkDestroyDevice(vk_device, NULL); + vkDestroyInstance(instance, NULL); end: if (error) cl_util_print_error(error); return error; diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index 6bf6de0a..306c9c24 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -231,9 +231,6 @@ int main(int argc, char* argv[]) << std::endl; } - // Create OpenCL runtime objects. - cl::Context cl_context{ cl_device }; - // Check if the device supports the Khronos extensions needed before // attempting to compile the kernel. if (diag_opts.verbose) @@ -249,10 +246,15 @@ int main(int argc, char* argv[]) std::cout << "OpenCL device does not support the required " "Khronos extension " << extension << std::endl; + vkDestroyDevice(vk_device, nullptr); + vkDestroyInstance(instance, nullptr); exit(EXIT_SUCCESS); } } + // Create OpenCL runtime objects. + cl::Context cl_context{ cl_device }; + // Compile kernel. if (diag_opts.verbose) { @@ -596,6 +598,8 @@ int main(int argc, char* argv[]) vkUnmapMemory(vk_device, vk_buf_x_memory); vkFreeMemory(vk_device, vk_buf_y_memory, nullptr); vkFreeMemory(vk_device, vk_buf_x_memory, nullptr); + vkDestroyDevice(vk_device, nullptr); + vkDestroyInstance(instance, nullptr); } catch (cl::BuildError& e) { From 72fc5af688d42021d03720ec2cd7aed03d7e2eb3 Mon Sep 17 00:00:00 2001 From: Xin Jin Date: Wed, 30 Jul 2025 15:11:37 +0100 Subject: [PATCH 2/4] Fix use-after-free on dev_version in version logic Fix a use-after-free warning reported by GCC: error: pointer 'dev_version' may be used after 'free' else if (opencl_version_contains(dev_version, "2.")) Previously, dev_version was freed immediately after checking for OpenCL 1.0/1.1, but later reused for determining whether to add -cl-std compiler options. This triggered -Werror=use-after-free under GCC with strict diagnostics. This patch: - Defers free(dev_version) to a new 'ver:' cleanup label to avoid premature free. - Replaces raw malloc/free with MEM_CHECK for safety and consistency. - Adds snprintf-based -cl-std option handling with bounds checking. - Moves compiler_options logic earlier to avoid stale pointer usage. - Removes old compiler_options block near clBuildProgram. The refactoring ensures correctness, eliminates UB, and prepares for cleaner version-based behavior in the future. Signed-off-by: Xin Jin --- samples/core/multi-device/main.c | 73 ++++++++++++++++++++------------ 1 file changed, 45 insertions(+), 28 deletions(-) diff --git a/samples/core/multi-device/main.c b/samples/core/multi-device/main.c index 2bf2a856..a5061466 100644 --- a/samples/core/multi-device/main.c +++ b/samples/core/multi-device/main.c @@ -221,16 +221,18 @@ int main(int argc, char* argv[]) // Query OpenCL version supported by device. size_t dev_version_size; - OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_VERSION, 0, NULL, - &dev_version_size), - error, end); - char *dev_version = malloc(dev_version_size); - if (!dev_version) - { - fprintf(stderr, "Failed to allocate memory for device version.\n"); - exit(EXIT_FAILURE); - } + OCLERROR_RET( + clGetDeviceInfo(dev, CL_DEVICE_VERSION, 0, NULL, &dev_version_size), + error, end); + + char compiler_options[1023] = ""; + char* dev_version = NULL; + MEM_CHECK(dev_version = (char*)malloc(dev_version_size), error, end); + + OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_VERSION, dev_version_size, + dev_version, NULL), + error, ver); if (opencl_version_contains(dev_version, "1.0") || opencl_version_contains(dev_version, "1.1")) @@ -240,17 +242,44 @@ int main(int argc, char* argv[]) "1.2 feature, but the device chosen only supports OpenCL %s. " "Please try with a different OpenCL device instead.\n", dev_version); - free(dev_version); - exit(EXIT_SUCCESS); + error = CL_SUCCESS; + goto ver; + } + else + { + // If no -cl-std option is specified then the highest 1.x version + // supported by each device is used to compile the program. Therefore, + // it's only necessary to add the -cl-std option for 2.0 and 3.0 OpenCL + // versions. + + int written = 0; + if (opencl_version_contains(dev_version, "3.")) + { + written = snprintf(compiler_options, sizeof(compiler_options), + "-cl-std=CL3.0 "); + } + else if (opencl_version_contains(dev_version, "2.")) + { + written = snprintf(compiler_options, sizeof(compiler_options), + "-cl-std=CL2.0 "); + } + + if (written < 0 || written >= (int)sizeof(compiler_options)) + { + fprintf( + stderr, + "Error: compiler_options buffer overflow or encoding error.\n"); + free(dev_version); + exit(EXIT_FAILURE); + } } - free(dev_version); // Check if device supports fission. cl_device_partition_property* dev_props = NULL; size_t props_size = 0; OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_PARTITION_PROPERTIES, 0, NULL, &props_size), - error, end); + error, ver); if (props_size == 0) { fprintf(stdout, @@ -263,7 +292,7 @@ int main(int argc, char* argv[]) // Check if the "partition equally" type is supported. MEM_CHECK(dev_props = (cl_device_partition_property*)malloc(props_size), - error, end); + error, ver); OCLERROR_RET(clGetDeviceInfo(dev, CL_DEVICE_PARTITION_PROPERTIES, props_size, dev_props, NULL), error, props); @@ -337,20 +366,6 @@ int main(int argc, char* argv[]) context, 1, (const char**)&kernel, &program_size, &error), error, ker); - // If no -cl-std option is specified then the highest 1.x version - // supported by each device is used to compile the program. Therefore, - // it's only necessary to add the -cl-std option for 2.0 and 3.0 OpenCL - // versions. - char compiler_options[1023] = ""; - if (opencl_version_contains(dev_version, "3.")) - { - strcat(compiler_options, "-cl-std=CL3.0 "); - } - else if (opencl_version_contains(dev_version, "2.")) - { - strcat(compiler_options, "-cl-std=CL2.0 "); - } - OCLERROR_RET( clBuildProgram(program, 2, subdevices, compiler_options, NULL, NULL), error, prg); @@ -710,6 +725,8 @@ int main(int argc, char* argv[]) free(subdevices); props: free(dev_props); +ver: + if (dev_version) free(dev_version); end: if (error) cl_util_print_error(error); return error; From a27b74ea2ffa08788c3da434662188edf1691cc8 Mon Sep 17 00:00:00 2001 From: Xin Jin Date: Wed, 30 Jul 2025 15:22:44 +0100 Subject: [PATCH 3/4] Replace label endl with end To address review comment Signed-off-by: Xin Jin --- samples/core/blur/main.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/samples/core/blur/main.c b/samples/core/blur/main.c index 025af91e..53c89b8a 100644 --- a/samples/core/blur/main.c +++ b/samples/core/blur/main.c @@ -344,7 +344,7 @@ cl_int single_pass_box_blur(state *const s, cl_int size) // compile kernel cl_kernel blur; OCLERROR_PAR(blur = clCreateKernel(s->program, "blur_box", &error), error, - endl); + end); // set kernel parameters OCLERROR_RET(clSetKernelArg(blur, 0, sizeof(cl_mem), &s->input_image_buf), @@ -375,7 +375,7 @@ cl_int single_pass_box_blur(state *const s, cl_int size) clReleaseEvent(pass); blr: clReleaseKernel(blur); -endl: +end: return error; } From 9cdff1aef3b35bd92966e8320cfb1c2071291515 Mon Sep 17 00:00:00 2001 From: Xin Jin Date: Thu, 7 Aug 2025 16:01:22 +0100 Subject: [PATCH 4/4] Remove redundant NULL check before free(dev_version) Calling free() on a NULL pointer is well-defined and has no effect. This simplifies the code by removing an unnecessary conditional. Signed-off-by: Xin Jin --- samples/core/multi-device/main.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/core/multi-device/main.c b/samples/core/multi-device/main.c index a5061466..96bdbf6e 100644 --- a/samples/core/multi-device/main.c +++ b/samples/core/multi-device/main.c @@ -726,7 +726,7 @@ int main(int argc, char* argv[]) props: free(dev_props); ver: - if (dev_version) free(dev_version); + free(dev_version); end: if (error) cl_util_print_error(error); return error;