From 094cc04e16a4f825b795999bf4575c535314bfb5 Mon Sep 17 00:00:00 2001 From: John Kesapides <46718829+JohnKesapidesARM@users.noreply.github.com> Date: Fri, 22 May 2020 13:26:05 +0100 Subject: [PATCH] Mem-leaks from conformance pipes (#772) Fix various memory leaks around events. Convert test to use supplied typewrappers to avoid memory leaks. Also use error helper functions to reduce code size. Use stringstreams to synthesize kernel sources, and raw c+11 string literals. Signed-off-by: John Kesapides --- test_conformance/pipes/test_pipe_info.cpp | 57 +- test_conformance/pipes/test_pipe_limits.cpp | 1131 +++++------------ .../pipes/test_pipe_query_functions.cpp | 429 +------ .../pipes/test_pipe_read_write.cpp | 647 ++++------ .../pipes/test_pipe_readwrite_errors.cpp | 234 +--- .../pipes/test_pipe_subgroups.cpp | 229 +--- 6 files changed, 790 insertions(+), 1937 deletions(-) diff --git a/test_conformance/pipes/test_pipe_info.cpp b/test_conformance/pipes/test_pipe_info.cpp index 5d3e3a49..7543c6cd 100644 --- a/test_conformance/pipes/test_pipe_info.cpp +++ b/test_conformance/pipes/test_pipe_info.cpp @@ -21,29 +21,25 @@ const char* pipe_kernel_code = { int test_pipe_info( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements ) { - cl_mem pipe; + clMemWrapper pipe; cl_int err; cl_uint pipe_width = 512; cl_uint pipe_depth = 1024; cl_uint returnVal; - cl_program program; - cl_kernel kernel; + clProgramWrapper program; + clKernelWrapper kernel; - pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, pipe_width, pipe_depth, NULL, &err); + pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, pipe_width, pipe_depth, + NULL, &err); test_error(err, "clCreatePipe failed."); - err = clGetPipeInfo(pipe, CL_PIPE_PACKET_SIZE, sizeof(pipe_width), (void *)&returnVal, NULL); - if ( err ) - { - log_error( "Error calling clGetPipeInfo(): %d\n", err ); - clReleaseMemObject(pipe); - return -1; - } + err = clGetPipeInfo(pipe, CL_PIPE_PACKET_SIZE, sizeof(pipe_width), + (void *)&returnVal, NULL); + test_error(err, "clGetPipeInfo failed."); - if(pipe_width != returnVal) + if (pipe_width != returnVal) { - log_error( "Error in clGetPipeInfo() check of pipe packet size\n" ); - clReleaseMemObject(pipe); + log_error("Error in clGetPipeInfo() check of pipe packet size\n"); return -1; } else @@ -52,17 +48,11 @@ int test_pipe_info( cl_device_id deviceID, cl_context context, cl_command_queue } err = clGetPipeInfo(pipe, CL_PIPE_MAX_PACKETS, sizeof(pipe_depth), (void *)&returnVal, NULL); - if ( err ) - { - log_error( "Error calling clGetPipeInfo(): %d\n", err ); - clReleaseMemObject(pipe); - return -1; - } + test_error(err, "clGetPipeInfo failed."); if(pipe_depth != returnVal) { log_error( "Error in clGetPipeInfo() check of pipe max packets\n" ); - clReleaseMemObject(pipe); return -1; } else @@ -71,39 +61,20 @@ int test_pipe_info( cl_device_id deviceID, cl_context context, cl_command_queue } err = create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, (const char**)&pipe_kernel_code, "pipe_kernel", "-cl-std=CL2.0 -cl-kernel-arg-info"); - if(err) - { - clReleaseMemObject(pipe); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " Error creating program", -1); cl_kernel_arg_type_qualifier arg_type_qualifier = 0; cl_kernel_arg_type_qualifier expected_type_qualifier = CL_KERNEL_ARG_TYPE_PIPE; err = clGetKernelArgInfo( kernel, 0, CL_KERNEL_ARG_TYPE_QUALIFIER, sizeof(arg_type_qualifier), &arg_type_qualifier, NULL ); - if(err) - { - clReleaseMemObject(pipe); - clReleaseKernel(kernel); - clReleaseProgram(program); - print_error(err, "clSetKernelArg failed\n"); - return -1; - } + test_error_ret(err, " clSetKernelArgInfo failed", -1); err = (arg_type_qualifier != expected_type_qualifier); + if(err) { - clReleaseMemObject(pipe); - clReleaseKernel(kernel); - clReleaseProgram(program); print_error(err, "ERROR: Bad type qualifier\n"); return -1; } - // cleanup - clReleaseMemObject(pipe); - clReleaseKernel(kernel); - clReleaseProgram(program); - return err; } diff --git a/test_conformance/pipes/test_pipe_limits.cpp b/test_conformance/pipes/test_pipe_limits.cpp index 29199df6..85247f82 100644 --- a/test_conformance/pipes/test_pipe_limits.cpp +++ b/test_conformance/pipes/test_pipe_limits.cpp @@ -15,63 +15,103 @@ // #include "harness/compat.h" +#include +#include +#include +#include #include #include -#include -#include +#include #include +#include #include "procs.h" #include "harness/errorHelpers.h" #define STRING_LENGTH 1024 -void createKernelSourceCode(char *source, int num_pipes) +void createKernelSourceCode(std::stringstream &stream, int num_pipes) { int i; - char str[256]; - int str_length; - strcpy(source, "__kernel void test_multiple_pipe_write(__global int *src, "); + stream << "__kernel void test_multiple_pipe_write(__global int *src, "; + for (i = 0; i < num_pipes; i++) + { + stream << "__write_only pipe int pipe" << i << ", "; + } + stream << R"(int num_pipes ) + { + int gid = get_global_id(0); + reserve_id_t res_id; - for(i = 0; i < num_pipes; i++) { - sprintf(str, "__write_only pipe int pipe%d, ", i); - strcat(source, str); - } - sprintf(str, "int num_pipes ) \n{\n int gid = get_global_id(0);\n reserve_id_t res_id;\n\n"); - strcat(source, str); - sprintf(str, " if(gid < (get_global_size(0))/num_pipes)\n {\n res_id = reserve_write_pipe(pipe0, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " write_pipe(pipe0, res_id, 0, &src[gid]);\n commit_write_pipe(pipe0, res_id);\n }\n }\n"); - strcat(source, str); - for(i = 1; i < num_pipes; i++){ - sprintf(str, " else if(gid < (%d*get_global_size(0))/num_pipes)\n {\n res_id = reserve_write_pipe(pipe%d, 1);\n if(is_valid_reserve_id(res_id))\n {\n", i+1, i); - strcat(source, str); - sprintf(str, " write_pipe(pipe%d, res_id, 0, &src[gid]);\n commit_write_pipe(pipe%d, res_id);\n }\n }\n", i, i); - strcat(source, str); - } - strcat(source, "}\n\n__kernel void test_multiple_pipe_read(__global int *dst, "); - for(i = 0; i < num_pipes; i++) { - sprintf(str, "__read_only pipe int pipe%d, ", i); - strcat(source, str); - } - sprintf(str, "int num_pipes ) \n{\n int gid = get_global_id(0);\n reserve_id_t res_id;\n\n"); - strcat(source, str); - sprintf(str, " if(gid < (get_global_size(0))/num_pipes)\n {\n res_id = reserve_read_pipe(pipe0, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " read_pipe(pipe0, res_id, 0, &dst[gid]);\n commit_read_pipe(pipe0, res_id);\n }\n }\n"); - strcat(source, str); - for(i = 1; i < num_pipes; i++){ - sprintf(str, " else if(gid < (%d*get_global_size(0))/num_pipes)\n {\n res_id = reserve_read_pipe(pipe%d, 1);\n if(is_valid_reserve_id(res_id))\n {\n", i+1, i); - strcat(source, str); - sprintf(str, " read_pipe(pipe%d, res_id, 0, &dst[gid]);\n commit_read_pipe(pipe%d, res_id);\n }\n }\n", i, i); - strcat(source, str); - } - strcat(source, "}"); + if(gid < (get_global_size(0))/num_pipes) + { + res_id = reserve_write_pipe(pipe0, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(pipe0, res_id, 0, &src[gid]); + commit_write_pipe(pipe0, res_id); + } + })"; - str_length = strlen(source); - assert(str_length <= STRING_LENGTH*num_pipes); + for (i = 1; i < num_pipes; i++) + { + // clang-format off + stream << R"( + else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes) + { + res_id = reserve_write_pipe(pipe)" << i << R"(, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(pipe)" << i << R"(, res_id, 0, &src[gid]); + commit_write_pipe(pipe)" << i << R"(, res_id); + } + } + )"; + // clang-format om + } + stream << R"( + } + + __kernel void test_multiple_pipe_read(__global int *dst, )"; + + for (i = 0; i < num_pipes; i++) + { + stream << "__read_only pipe int pipe" << i << ", "; + } + stream << R"(int num_pipes ) + { + int gid = get_global_id(0); + reserve_id_t res_id; + + + if(gid < (get_global_size(0))/num_pipes) + { + res_id = reserve_read_pipe(pipe0, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(pipe0, res_id, 0, &dst[gid]); + commit_read_pipe(pipe0, res_id); + } + })"; + + for (i = 1; i < num_pipes; i++) + { + // clang-format off + stream << R"( + else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes) + { + res_id = reserve_read_pipe(pipe)" << i << R"(, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(pipe)" << i << R"(, res_id, 0, &dst[gid]); + commit_read_pipe(pipe)" << i << R"(, res_id); + } + })"; + // clang-format on + } + stream << "}"; } static int verify_result(void *ptr1, void *ptr2, int n) @@ -113,28 +153,34 @@ static int verify_result_int(void *ptr1, void *ptr2, int n) int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipes[1024]; - cl_mem buffers[2]; - void *outptr; - cl_int *inptr; - cl_program program; - cl_kernel kernel[2]; - size_t global_work_size[3]; - cl_int err; - cl_int size; - int num_pipe_elements = 1024; - int i, j; - int max_pipe_args; - char *source; - cl_event producer_sync_event = NULL; - cl_event consumer_sync_event = NULL; - MTdata d = init_genrand( gRandomSeed ); - const char* kernelName[] = {"test_multiple_pipe_write", "test_multiple_pipe_read"}; + clMemWrapper pipes[1024]; + clMemWrapper buffers[2]; + void *outptr; + cl_int *inptr; + clProgramWrapper program; + clKernelWrapper kernel[2]; + size_t global_work_size[3]; + cl_int err; + cl_int size; + int num_pipe_elements = 1024; + int i, j; + int max_pipe_args; + std::stringstream source; + clEventWrapper producer_sync_event = NULL; + clEventWrapper consumer_sync_event = NULL; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr; - size_t min_alignment = get_min_alignment(context); + MTdataHolder d(gRandomSeed); + const char *kernelName[] = { "test_multiple_pipe_write", + "test_multiple_pipe_read" }; - err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, sizeof(max_pipe_args), (void*)&max_pipe_args, NULL); - if(err){ + size_t min_alignment = get_min_alignment(context); + + err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, + sizeof(max_pipe_args), (void *)&max_pipe_args, NULL); + if (err) + { print_error(err, " clGetDeviceInfo failed\n"); return -1; } @@ -145,76 +191,41 @@ int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_que global_work_size[0] = (cl_uint)num_pipe_elements * max_pipe_args; size = sizeof(int) * num_pipe_elements * max_pipe_args; - source = (char *)malloc(STRING_LENGTH * sizeof(char) * max_pipe_args); inptr = (cl_int *)align_malloc(size, min_alignment); for(i = 0; i < num_pipe_elements * max_pipe_args; i++){ inptr[i] = (int)genrand_int32(d); } + BufferInPtr.reset(inptr, nullptr, 0, size, true); buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - free(source); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + outptr = align_malloc(size, min_alignment); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - free(source); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); for(i = 0; i < max_pipe_args; i++){ pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_pipe_elements, NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - free(source); - for(j = 0; j < i; j++) { - clReleaseMemObject(pipes[j]); - } - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); } createKernelSourceCode(source, max_pipe_args); + std::string kernel_source = source.str(); + const char *sources[] = { kernel_source.c_str() }; + // Create producer kernel - err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - align_free(outptr); - free(source); - print_error(err, "Error creating program\n"); - return -1; - } + err = create_single_kernel_helper_with_build_options( + context, &program, &kernel[0], 1, sources, kernelName[0], + "-cl-std=CL2.0"); + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - align_free(outptr); - free(source); - print_error(err, " Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); for( i = 0; i < max_pipe_args; i++){ @@ -226,127 +237,28 @@ int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_que err |= clSetKernelArg(kernel[1], i+1, sizeof(cl_mem), (void*)&pipes[i]); } err |= clSetKernelArg(kernel[1], max_pipe_args + 1, sizeof(int), (void*)&max_pipe_args); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - free(source); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); // Launch Consumer kernel err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clWaitForEvents(1, &consumer_sync_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clWaitForEvents failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clWaitForEvents failed", -1); if( verify_result( inptr, outptr, num_pipe_elements*sizeof(cl_int))){ log_error("test_pipe_max_args failed\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; } else { log_info("test_pipe_max_args passed\n"); } - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - for(j = 0; j < max_pipe_args; j++) { - clReleaseMemObject(pipes[j]); - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); return 0; } @@ -354,39 +266,41 @@ int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_que int test_pipe_max_packet_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipe; - cl_mem buffers[2]; - void *outptr; - cl_char *inptr; - cl_program program; - cl_kernel kernel[2]; - size_t global_work_size[3]; - cl_int err; - size_t size; - int num_pipe_elements = 1024; - int i; - cl_uint max_pipe_packet_size; - char *source; - char str[256]; - int str_length; - cl_event producer_sync_event = NULL; - cl_event consumer_sync_event = NULL; - MTdata d = init_genrand( gRandomSeed ); - const char* kernelName[] = {"test_pipe_max_packet_size_write", "test_pipe_max_packet_size_read"}; + clMemWrapper pipe; + clMemWrapper buffers[2]; + void *outptr; + cl_char *inptr; + clProgramWrapper program; + clKernelWrapper kernel[2]; + size_t global_work_size[3]; + cl_int err; + size_t size; + int num_pipe_elements = 1024; + int i; + cl_uint max_pipe_packet_size; + clEventWrapper producer_sync_event = NULL; + clEventWrapper consumer_sync_event = NULL; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr; + MTdataHolder d(gRandomSeed); + const char *kernelName[] = { "test_pipe_max_packet_size_write", + "test_pipe_max_packet_size_read" }; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); global_work_size[0] = (cl_uint)num_pipe_elements; - source = (char*)malloc(STRING_LENGTH*sizeof(char)); + std::stringstream source; - err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE, sizeof(max_pipe_packet_size), (void*)&max_pipe_packet_size, NULL); - if(err){ - print_error(err, " clGetDeviceInfo failed\n"); - return -1; - } - if(max_pipe_packet_size < 1024){ - log_error("The device should support minimum packet size of 1024 bytes"); + err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE, + sizeof(max_pipe_packet_size), + (void *)&max_pipe_packet_size, NULL); + test_error_ret(err, " clCreatePipe failed", -1); + + if (max_pipe_packet_size < 1024) + { + log_error( + "The device should support minimum packet size of 1024 bytes"); return -1; } @@ -402,354 +316,248 @@ int test_pipe_max_packet_size(cl_device_id deviceID, cl_context context, cl_comm for(i = 0; i < size; i++){ inptr[i] = (char)genrand_int32(d); } + BufferInPtr.reset(inptr, nullptr, 0, size, true); buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - free(source); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + outptr = align_malloc(size, min_alignment); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); + buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - free(source); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, max_pipe_packet_size, num_pipe_elements, NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - free(source); - clReleaseMemObject(pipe); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); - sprintf(str, "typedef struct{\n char a[%d];\n}TestStruct;\n\n__kernel void test_pipe_max_packet_size_write(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n{\n", max_pipe_packet_size); - strcpy(source,str); - strcat(source, " int gid = get_global_id(0);\n reserve_id_t res_id;\n\n"); - sprintf(str, " res_id = reserve_write_pipe(out_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " write_pipe(out_pipe, res_id, 0, &src[gid]);\n commit_write_pipe(out_pipe, res_id);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_max_packet_size_read(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n{\n"); - strcat(source, str); - strcat(source, " int gid = get_global_id(0);\n reserve_id_t res_id;\n\n"); - sprintf(str, " res_id = reserve_read_pipe(in_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n commit_read_pipe(in_pipe, res_id);\n }\n}\n\n"); - strcat(source, str); + // clang-format off + source << R"( + typedef struct{ + char a[)" << max_pipe_packet_size << R"(]; + }TestStruct; - str_length = strlen(source); - assert(str_length <= STRING_LENGTH); + __kernel void test_pipe_max_packet_size_write(__global TestStruct *src, __write_only pipe TestStruct out_pipe) + { + int gid = get_global_id(0); + reserve_id_t res_id; + + res_id = reserve_write_pipe(out_pipe, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(out_pipe, res_id, 0, &src[gid]); + commit_write_pipe(out_pipe, res_id); + } + } + + __kernel void test_pipe_max_packet_size_read(__read_only pipe TestStruct in_pipe, __global TestStruct *dst) + { + int gid = get_global_id(0); + reserve_id_t res_id; + + res_id = reserve_read_pipe(in_pipe, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(in_pipe, res_id, 0, &dst[gid]); + commit_read_pipe(in_pipe, res_id); + } + } + )"; + // clang-format on + + std::string kernel_source = source.str(); + const char *sources[] = { kernel_source.c_str() }; // Create producer kernel - err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - align_free(outptr); - free(source); - print_error(err, "Error creating program\n"); - return -1; - } + err = create_single_kernel_helper_with_build_options( + context, &program, &kernel[0], 1, sources, kernelName[0], + "-cl-std=CL2.0"); + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - align_free(outptr); - free(source); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - free(source); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); // Launch Consumer kernel err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if( verify_result( inptr, outptr, size)){ log_error("test_pipe_max_packet_size failed\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); - return -1; } else { log_info("test_pipe_max_packet_size passed\n"); } - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - free(source); return 0; } int test_pipe_max_active_reservations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipe; - cl_mem buffers[2]; - cl_mem buf_reservations; - cl_mem buf_status; - cl_mem buf_reserve_id_t_size; - cl_mem buf_reserve_id_t_size_aligned; - cl_int *inptr; - void *outptr; - int size, i; - cl_program program; - cl_kernel kernel[3]; - size_t global_work_size[3]; - cl_int err; - int status = 0; - cl_uint max_active_reservations = 0; - cl_ulong max_global_size = 0; - int reserve_id_t_size; - int temp; - char *source; - char str[256]; - int str_length; - cl_event sync_event = NULL; - cl_event read_event = NULL; - MTdata d = init_genrand( gRandomSeed ); - const char* kernelName[3] = {"test_pipe_max_active_reservations_write", "test_pipe_max_active_reservations_read", "pipe_get_reserve_id_t_size"}; + clMemWrapper pipe; + clMemWrapper buffers[2]; + clMemWrapper buf_reservations; + clMemWrapper buf_status; + clMemWrapper buf_reserve_id_t_size; + clMemWrapper buf_reserve_id_t_size_aligned; + cl_int *inptr; + void *outptr; + int size, i; + clProgramWrapper program; + clKernelWrapper kernel[3]; + size_t global_work_size[3]; + cl_int err; + int status = 0; + cl_uint max_active_reservations = 0; + cl_ulong max_global_size = 0; + int reserve_id_t_size; + int temp; + clEventWrapper sync_event = NULL; + clEventWrapper read_event = NULL; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr; + MTdataHolder d(gRandomSeed); + const char *kernelName[3] = { "test_pipe_max_active_reservations_write", + "test_pipe_max_active_reservations_read", + "pipe_get_reserve_id_t_size" }; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); - source = (char*)malloc(2*STRING_LENGTH*sizeof(char)); + std::stringstream source; global_work_size[0] = 1; - err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, sizeof(max_active_reservations), (void*)&max_active_reservations, NULL); - if(err){ - print_error(err, " clGetDeviceInfo failed\n"); - return -1; - } + err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, + sizeof(max_active_reservations), + (void *)&max_active_reservations, NULL); + test_error_ret(err, " clGetDeviceInfo failed", -1); - err = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(max_global_size), (void*)&max_global_size, NULL); - if(err){ - print_error(err, " clGetDeviceInfo failed\n"); - return -1; - } + err = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(max_global_size), (void *)&max_global_size, + NULL); + test_error_ret(err, " clGetDeviceInfo failed", -1); - max_active_reservations = (max_active_reservations > max_global_size) ? 1<<16 : max_active_reservations; + max_active_reservations = (max_active_reservations > max_global_size) + ? 1 << 16 + : max_active_reservations; - if(max_active_reservations < 1){ + if (max_active_reservations < 1) + { log_error("The device should support minimum active reservations of 1"); return -1; } // To get reserve_id_t size buf_reserve_id_t_size = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, sizeof(reserve_id_t_size), NULL, &err); - if ( err ){ - clReleaseMemObject(buf_reserve_id_t_size); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); - sprintf(str, "__kernel void test_pipe_max_active_reservations_write(__global int *src, __write_only pipe int out_pipe, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)\n{\n"); - strcpy(source,str); - sprintf(str, " __global reserve_id_t *res_id_ptr;\n int reserve_idx;\n int commit_idx;\n"); - strcat(source, str); - sprintf(str, " for(reserve_idx = 0; reserve_idx < %d; reserve_idx++)\n {\n", max_active_reservations); - strcat(source, str); - sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);\n"); - strcat(source, str); - sprintf(str, " *res_id_ptr = reserve_write_pipe(out_pipe, 1);\n"); - strcat(source, str); - sprintf(str, " if(is_valid_reserve_id(res_id_ptr[0]))\n {\n write_pipe(out_pipe, res_id_ptr[0], 0, &src[reserve_idx]);\n }\n"); - strcat(source, str); - sprintf(str, " else\n {\n *status = -1;\n return;\n }\n }\n"); - strcat(source, str); - sprintf(str, " for(commit_idx = 0; commit_idx < %d; commit_idx++)\n {\n", max_active_reservations); - strcat(source, str); - sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);\n"); - strcat(source, str); - sprintf(str, " commit_write_pipe(out_pipe, res_id_ptr[0]);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_max_active_reservations_read(__read_only pipe int in_pipe, __global int *dst, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)\n{\n"); - strcat(source, str); - sprintf(str, " __global reserve_id_t *res_id_ptr;\n int reserve_idx;\n int commit_idx;\n"); - strcat(source, str); - sprintf(str, " for(reserve_idx = 0; reserve_idx < %d; reserve_idx++)\n {\n", max_active_reservations); - strcat(source, str); - sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);\n"); - strcat(source, str); - sprintf(str, " *res_id_ptr = reserve_read_pipe(in_pipe, 1);\n"); - strcat(source, str); - sprintf(str, " if(is_valid_reserve_id(res_id_ptr[0]))\n {\n read_pipe(in_pipe, res_id_ptr[0], 0, &dst[reserve_idx]);\n }\n"); - strcat(source, str); - sprintf(str, " else\n {\n *status = -1;\n return;\n }\n }\n"); - strcat(source, str); - sprintf(str, " for(commit_idx = 0; commit_idx < %d; commit_idx++)\n {\n", max_active_reservations); - strcat(source, str); - sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);\n"); - strcat(source, str); - sprintf(str, " commit_read_pipe(in_pipe, res_id_ptr[0]);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void pipe_get_reserve_id_t_size(__global int *reserve_id_t_size) \n"); - strcat(source, str); - sprintf(str, "{\n *reserve_id_t_size = sizeof(reserve_id_t);\n}\n"); - strcat(source, str); + // clang-format off + source << R"( + __kernel void test_pipe_max_active_reservations_write(__global int *src, __write_only pipe int out_pipe, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status) + { + __global reserve_id_t *res_id_ptr; + int reserve_idx; + int commit_idx; - str_length = strlen(source); - assert(str_length <= 2*STRING_LENGTH); + for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++) + { + res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]); + *res_id_ptr = reserve_write_pipe(out_pipe, 1); + if(is_valid_reserve_id(res_id_ptr[0])) + { + write_pipe(out_pipe, res_id_ptr[0], 0, &src[reserve_idx]); + } + else + { + *status = -1; + return; + } + } + + for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++) + { + res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]); + commit_write_pipe(out_pipe, res_id_ptr[0]); + } + } + + __kernel void test_pipe_max_active_reservations_read(__read_only pipe int in_pipe, __global int *dst, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status) + { + __global reserve_id_t *res_id_ptr; + int reserve_idx; + int commit_idx; + + for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++) + { + res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]); + *res_id_ptr = reserve_read_pipe(in_pipe, 1); + + if(is_valid_reserve_id(res_id_ptr[0])) + { + read_pipe(in_pipe, res_id_ptr[0], 0, &dst[reserve_idx]); + } + else + { + *status = -1; + return; + } + } + + for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++) + { + res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]); + commit_read_pipe(in_pipe, res_id_ptr[0]); + } + } + + __kernel void pipe_get_reserve_id_t_size(__global int *reserve_id_t_size) + { + *reserve_id_t_size = sizeof(reserve_id_t); + } + )"; + // clang-format on + + std::string kernel_source = source.str(); + const char *sources[] = { kernel_source.c_str() }; // Create producer kernel - err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buf_reserve_id_t_size); - print_error(err, "Error creating program\n"); - return -1; - } + err = create_single_kernel_helper_with_build_options( + context, &program, &kernel[0], 1, sources, kernelName[0], + "-cl-std=CL2.0"); + test_error_ret(err, " Error creating program", -1); // Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buf_reserve_id_t_size); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); // Create size query kernel for reserve_id_t kernel[2] = clCreateKernel(program, kernelName[2], &err); - if( kernel[2] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buf_reserve_id_t_size); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); + err = clSetKernelArg(kernel[2], 0, sizeof(cl_mem), (void*)&buf_reserve_id_t_size); - if(err){ - clReleaseMemObject(buf_reserve_id_t_size); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); + //Launch size query kernel for reserve_id_t err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buf_reserve_id_t_size, true, 0, sizeof(reserve_id_t_size), &reserve_id_t_size, 1, &sync_event, &read_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); err = clWaitForEvents(1, &read_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clWaitForEvents failed" ); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clWaitForEvents failed", -1); // Round reserve_id_t_size to the nearest power of 2 temp = 1; @@ -763,323 +571,84 @@ int test_pipe_max_active_reservations(cl_device_id deviceID, cl_context context, for(i = 0; i < max_active_reservations; i++){ inptr[i] = (int)genrand_int32(d); } + BufferInPtr.reset(inptr, nullptr, 0, size, true); buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if ( err ){ - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buffers[0]); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); outptr = align_malloc(size, min_alignment); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); + buffers[1] = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, size, NULL, &err); - if ( err ){ - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - align_free(outptr); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); buf_reserve_id_t_size_aligned = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(reserve_id_t_size), &reserve_id_t_size, &err); - if ( err ){ - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); //For error status buf_status = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err); - if ( err ){ - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), max_active_reservations, NULL, &err); - if(err){ - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); // Global buffer to hold all active reservation ids buf_reservations = clCreateBuffer(context, CL_MEM_HOST_NO_ACCESS, reserve_id_t_size*max_active_reservations, NULL, &err); - if ( err != CL_SUCCESS ){ - print_error( err, " clCreateBuffer failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseEvent(read_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buf_reservations); err |= clSetKernelArg(kernel[0], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned); err |= clSetKernelArg(kernel[0], 4, sizeof(cl_mem), (void*)&buf_status); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseProgram(program); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); err = clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]); err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buf_reservations); err |= clSetKernelArg(kernel[1], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned); err |= clSetKernelArg(kernel[1], 4, sizeof(cl_mem), (void*)&buf_status); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseProgram(program); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); + + clReleaseEvent(sync_event); // Launch Producer kernel err = clEnqueueNDRangeKernel(queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if(status != 0) { log_error("test_pipe_max_active_reservations failed\n"); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(sync_event); - clReleaseProgram(program); return -1; } + clReleaseEvent(sync_event); // Launch Consumer kernel err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if(status != 0) { log_error("test_pipe_max_active_reservations failed\n"); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(sync_event); - clReleaseProgram(program); return -1; } err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if( verify_result_int( inptr, outptr, max_active_reservations)){ log_error("test_pipe_max_active_reservations failed\n"); - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(sync_event); - clReleaseProgram(program); return -1; } else { log_info("test_pipe_max_active_reservations passed\n"); } - //cleanup - clReleaseMemObject(buf_status); - clReleaseMemObject(buf_reserve_id_t_size); - clReleaseMemObject(buf_reserve_id_t_size_aligned); - clReleaseMemObject(buf_reservations); - clReleaseMemObject(pipe); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free(outptr); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(sync_event); - clReleaseProgram(program); return 0; } \ No newline at end of file diff --git a/test_conformance/pipes/test_pipe_query_functions.cpp b/test_conformance/pipes/test_pipe_query_functions.cpp index c1b4d922..f9c93aa2 100644 --- a/test_conformance/pipes/test_pipe_query_functions.cpp +++ b/test_conformance/pipes/test_pipe_query_functions.cpp @@ -79,28 +79,32 @@ static int verify_result(void *ptr1, void *ptr2, int n) int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipe; - cl_mem buffers[4]; - void *outptr1; - void *outptr2; - cl_int *inptr; - cl_program program; - cl_kernel kernel[3]; - size_t global_work_size[3]; - size_t half_global_work_size[3]; - size_t global_work_size_pipe_query[3]; - cl_int pipe_max_packets, pipe_num_packets; - cl_int err; - cl_int size; - cl_int i; - cl_event producer_sync_event = NULL; - cl_event consumer_sync_event = NULL; - cl_event pipe_query_sync_event = NULL; - cl_event pipe_read_sync_event = NULL; - MTdata d = init_genrand( gRandomSeed ); - const char* kernelName[] = {"test_pipe_write", "test_pipe_read", "test_pipe_query_functions"}; + clMemWrapper pipe; + clMemWrapper buffers[4]; + void *outptr1; + void *outptr2; + cl_int *inptr; + clProgramWrapper program; + clKernelWrapper kernel[3]; + size_t global_work_size[3]; + size_t half_global_work_size[3]; + size_t global_work_size_pipe_query[3]; + cl_int pipe_max_packets, pipe_num_packets; + cl_int err; + cl_int size; + cl_int i; + clEventWrapper producer_sync_event = NULL; + clEventWrapper consumer_sync_event = NULL; + clEventWrapper pipe_query_sync_event = NULL; + clEventWrapper pipe_read_sync_event = NULL; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr1; + BufferOwningPtr BufferOutPtr2; + MTdataHolder d(gRandomSeed); + const char *kernelName[] = { "test_pipe_write", "test_pipe_read", + "test_pipe_query_functions" }; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); size = sizeof(int) * num_elements; global_work_size[0] = (cl_uint)num_elements; @@ -109,98 +113,43 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm inptr = (int *)align_malloc(size, min_alignment); - for(i = 0; i < num_elements; i++){ + for (i = 0; i < num_elements; i++) + { inptr[i] = TEST_PRIME_INT; } + BufferInPtr.reset(inptr, nullptr, 0, size, true); buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + outptr1 = align_malloc(size/2, min_alignment); outptr2 = align_malloc(size, min_alignment); + BufferOutPtr1.reset(outptr1, nullptr, 0, size, true); + BufferOutPtr2.reset(outptr2, nullptr, 0, size, true); + buffers[1] = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, size, NULL, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr1 ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); buffers[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - align_free( outptr1 ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); buffers[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - align_free( outptr1 ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements, NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - align_free( outptr1 ); - clReleaseMemObject(pipe); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); // Create producer kernel err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_query_functions_kernel_code, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - align_free(outptr1); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - align_free(outptr1); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); + //Create pipe query functions kernel kernel[2] = clCreateKernel(program, kernelName[2], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - align_free(outptr1); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); @@ -209,104 +158,21 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm err |= clSetKernelArg(kernel[2], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[2], 1, sizeof(cl_mem), (void*)&buffers[2]); err |= clSetKernelArg(kernel[2], 2, sizeof(cl_mem), (void*)&buffers[3]); - - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseProgram(program); - align_free(outptr1); - print_error(err, " clSetKernelArg failed\n"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); // Launch Pipe query kernel err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &producer_sync_event, &pipe_query_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } - + test_error_ret(err, " clEnqueueReadBuffer failed", -1); err = clEnqueueReadBuffer(queue, buffers[3], true, 0, sizeof(cl_int), &pipe_max_packets, 1, &pipe_query_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if(pipe_num_packets != num_elements || pipe_max_packets != num_elements) { @@ -316,85 +182,20 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm // Launch Consumer kernel with half the previous global size err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, half_global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size / 2, outptr1, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); + + // We will reuse this variable so release the previous referred event. + clReleaseEvent(pipe_query_sync_event); // Launch Pipe query kernel err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &consumer_sync_event, &pipe_query_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, &pipe_read_sync_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); // After consumer kernel consumes num_elements/2 from the pipe, // there are (num_elements - num_elements/2) remaining package in the pipe. @@ -404,68 +205,24 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm return -1; } + // We will reuse this variable so release the previous referred event. + clReleaseEvent(producer_sync_event); + // Launch Producer kernel to fill the pipe again global_work_size[0] = pipe_num_packets; err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 1, &pipe_read_sync_event, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); + // We will reuse this variable so release the previous referred event. + clReleaseEvent(pipe_query_sync_event); // Launch Pipe query kernel err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &producer_sync_event, &pipe_query_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); + // We will reuse this variable so release the previous referred event. + clReleaseEvent(pipe_read_sync_event); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, &pipe_read_sync_event); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if(pipe_num_packets != num_elements) { @@ -473,48 +230,16 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm return -1; } + // We will reuse this variable so release the previous referred event. + clReleaseEvent(consumer_sync_event); + // Launch Consumer kernel to consume all packets from pipe global_work_size[0] = pipe_num_packets; err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &pipe_read_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr2, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed\n" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return -1; - } - + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if( verify_result(outptr1, outptr2, num_elements )){ log_error("test_pipe_query_functions failed\n"); @@ -523,22 +248,6 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm else { log_info("test_pipe_query_functions passed\n"); } - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(buffers[3]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseKernel(kernel[2]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseEvent(pipe_query_sync_event); - clReleaseEvent(pipe_read_sync_event); - clReleaseProgram(program); - align_free(outptr1); - return 0; } diff --git a/test_conformance/pipes/test_pipe_read_write.cpp b/test_conformance/pipes/test_pipe_read_write.cpp index 91a49909..4bb4468e 100644 --- a/test_conformance/pipes/test_pipe_read_write.cpp +++ b/test_conformance/pipes/test_pipe_read_write.cpp @@ -15,11 +15,15 @@ // #include "harness/compat.h" +#include +#include +#include +#include #include #include -#include +#include #include -#include +#include #include "procs.h" #include "kernels.h" @@ -89,113 +93,139 @@ static const char* convenience_float_kernel_name[] = { "test_pipe_convenience_wr static const char* convenience_half_kernel_name[] = { "test_pipe_convenience_write_half", "test_pipe_convenience_read_half", "test_pipe_convenience_write_half2", "test_pipe_convenience_read_half2", "test_pipe_convenience_write_half4", "test_pipe_convenience_read_half4", "test_pipe_convenience_write_half8", "test_pipe_convenience_read_half8", "test_pipe_convenience_write_half16", "test_pipe_convenience_read_half16" }; static const char* convenience_double_kernel_name[] = { "test_pipe_convenience_write_double", "test_pipe_convenience_read_double", "test_pipe_convenience_write_double2", "test_pipe_convenience_read_double2", "test_pipe_convenience_write_double4", "test_pipe_convenience_read_double4", "test_pipe_convenience_write_double8", "test_pipe_convenience_read_double8", "test_pipe_convenience_write_double16", "test_pipe_convenience_read_double16" }; -static void insertPragmaForHalfType(char *source, char *type) +static void insertPragmaForHalfType(std::stringstream &stream, char *type) { - source[0] = 0; - if(strncmp(type, "half",4) == 0) + if (strncmp(type, "half", 4) == 0) { - strcat(source, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"); + stream << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"; } } -void createKernelSource(char *source, char *type) +void createKernelSource(std::stringstream &stream, char *type) { - char str[512]; - int str_length; + insertPragmaForHalfType(stream, type); - insertPragmaForHalfType(source, type); + // clang-format off + stream << R"( + __kernel void test_pipe_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe) + { + int gid = get_global_id(0); + reserve_id_t res_id; - sprintf(str, "__kernel void test_pipe_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = reserve_write_pipe(out_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " write_pipe(out_pipe, res_id, 0, &src[gid]);\n commit_write_pipe(out_pipe, res_id);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = reserve_read_pipe(in_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n commit_read_pipe(in_pipe, res_id);\n }\n}\n"); - strcat(source, str); - str_length = strlen(source); - assert(str_length <= STRING_LENGTH); + res_id = reserve_write_pipe(out_pipe, 1); + if(is_valid_reserve_id(res_id)) + { + write_pipe(out_pipe, res_id, 0, &src[gid]); + commit_write_pipe(out_pipe, res_id); + } + } + + __kernel void test_pipe_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst) + { + int gid = get_global_id(0); + reserve_id_t res_id; + + res_id = reserve_read_pipe(in_pipe, 1); + if(is_valid_reserve_id(res_id)) + { + read_pipe(in_pipe, res_id, 0, &dst[gid]); + commit_read_pipe(in_pipe, res_id); + } + } + )"; + // clang-format on } -void createKernelSourceWorkGroup(char *source, char *type) +void createKernelSourceWorkGroup(std::stringstream &stream, char *type) { - char str[512]; - int str_length; + insertPragmaForHalfType(stream, type); - insertPragmaForHalfType(source, type); + // clang-format off + stream << R"( + __kernel void test_pipe_workgroup_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe) + { + int gid = get_global_id(0); + __local reserve_id_t res_id; - sprintf(str, "__kernel void test_pipe_workgroup_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n __local reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0));\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]);\n work_group_commit_write_pipe(out_pipe, res_id);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_workgroup_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n __local reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0));\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]);\n work_group_commit_read_pipe(in_pipe, res_id);\n }\n}\n"); - strcat(source, str); - str_length = strlen(source); - assert(str_length <= STRING_LENGTH); + res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]); + work_group_commit_write_pipe(out_pipe, res_id); + } + } + + __kernel void test_pipe_workgroup_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst) + { + int gid = get_global_id(0); + __local reserve_id_t res_id; + + res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0)); + if(is_valid_reserve_id(res_id)) + { + read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]); + work_group_commit_read_pipe(in_pipe, res_id); + } + } + )"; + // clang-format on } -void createKernelSourceSubGroup(char *source, char *type) +void createKernelSourceSubGroup(std::stringstream &stream, char *type) { - char str[512]; - int str_length; + insertPragmaForHalfType(stream, type); - insertPragmaForHalfType(source, type); + // clang-format off + stream << R"( + #pragma OPENCL EXTENSION cl_khr_subgroups : enable + __kernel void test_pipe_subgroup_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe) + { + int gid = get_global_id(0); + reserve_id_t res_id; - sprintf(str, "#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n__kernel void test_pipe_subgroup_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]);\n sub_group_commit_write_pipe(out_pipe, res_id);\n }\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_subgroup_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n"); - strcat(source, str); - sprintf(str, " res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());\n if(is_valid_reserve_id(res_id))\n {\n"); - strcat(source, str); - sprintf(str, " read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]);\n sub_group_commit_read_pipe(in_pipe, res_id);\n }\n}\n"); - strcat(source, str); - str_length = strlen(source); - assert(str_length <= STRING_LENGTH); + res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size()); + if(is_valid_reserve_id(res_id)) + { + write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]); + sub_group_commit_write_pipe(out_pipe, res_id); + } + } + + __kernel void test_pipe_subgroup_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst) + { + int gid = get_global_id(0); + reserve_id_t res_id; + + res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size()); + if(is_valid_reserve_id(res_id)) + { + read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]); + sub_group_commit_read_pipe(in_pipe, res_id); + } + } + )"; + // clang-format on } -void createKernelSourceConvenience(char *source, char *type) +void createKernelSourceConvenience(std::stringstream &stream, char *type) { - char str[512]; - int str_length; + insertPragmaForHalfType(stream, type); - insertPragmaForHalfType(source, type); + // clang-format off + stream << R"( + __kernel void test_pipe_convenience_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe) + { + int gid = get_global_id(0); + write_pipe(out_pipe, &src[gid]); + } - sprintf(str, "__kernel void test_pipe_convenience_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n write_pipe(out_pipe, &src[gid]);\n}\n\n"); - strcat(source, str); - sprintf(str, "__kernel void test_pipe_convenience_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type); - strcat(source, str); - sprintf(str, "{\n int gid = get_global_id(0);\n read_pipe(in_pipe, &dst[gid]);\n}\n"); - strcat(source, str); - str_length = strlen(source); - assert(str_length <= STRING_LENGTH); + __kernel void test_pipe_convenience_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst) + { + int gid = get_global_id(0); + read_pipe(in_pipe, &dst[gid]); + } + )"; + // clang-format on } // verify functions @@ -424,23 +454,24 @@ static int verify_readwrite_struct(void *ptr1, void *ptr2, int n) int test_pipe_readwrite( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops, void *inptr[5], const char *kernelName[], int (*fn)(void *, void *, int) ) { - cl_mem pipes[5]; - cl_mem buffers[10]; - void *outptr[5]; - cl_program program[5]; - cl_kernel kernel[10]; - size_t global_work_size[3]; - size_t local_work_size[3]; - cl_int err; - int i, ii; - size_t ptrSizes[5]; - int total_errors = 0; - cl_event producer_sync_event[5]; - cl_event consumer_sync_event[5]; - char *sourceCode[5]; - char vector_type[10]; + clMemWrapper pipes[5]; + clMemWrapper buffers[10]; + void *outptr[5]; + BufferOwningPtr BufferOutPtr[5]; + clProgramWrapper program[5]; + clKernelWrapper kernel[10]; + size_t global_work_size[3]; + size_t local_work_size[3]; + cl_int err; + int i, ii; + size_t ptrSizes[5]; + int total_errors = 0; + clEventWrapper producer_sync_event[5]; + clEventWrapper consumer_sync_event[5]; + std::stringstream sourceCode[5]; + char vector_type[10]; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); global_work_size[0] = (cl_uint)num_elements; @@ -450,217 +481,133 @@ int test_pipe_readwrite( cl_device_id deviceID, cl_context context, cl_command_q ptrSizes[3] = ptrSizes[2] << 1; ptrSizes[4] = ptrSizes[3] << 1; - for( i = 0; i < loops; i++) + for (i = 0; i < loops; i++) { ii = i << 1; - sourceCode[i] = (char*) malloc(STRING_LENGTH * sizeof(char)); - buffers[ii] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, ptrSizes[i] * num_elements, inptr[i], &err); - if(err){ - clReleaseMemObject(buffers[ii]); - align_free( outptr[i] ); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } - outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment); - buffers[ii+1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, ptrSizes[i] * num_elements, outptr[i], &err); - if ( err ){ - clReleaseMemObject(buffers[ii]); - align_free( outptr[i] ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + buffers[ii] = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + ptrSizes[i] * num_elements, inptr[i], &err); + test_error_ret(err, " clCreateBuffer failed", -1); + + outptr[i] = align_malloc(ptrSizes[i] * num_elements, min_alignment); + BufferOutPtr[i].reset(outptr[i], nullptr, 0, size, true); + buffers[ii + 1] = + clCreateBuffer(context, CL_MEM_USE_HOST_PTR, + ptrSizes[i] * num_elements, outptr[i], &err); + test_error_ret(err, " clCreateBuffer failed", -1); + // Creating pipe with non-power of 2 size - pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, ptrSizes[i], num_elements+3, NULL, &err); - if(err){ - clReleaseMemObject(pipes[i]); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, ptrSizes[i], + num_elements + 3, NULL, &err); + test_error_ret(err, " clCreatePipe failed", -1); - switch(i) + switch (i) { - case 0: - sprintf(vector_type, "%s", type); - break; - case 1: - sprintf(vector_type, "%s%d", type, 2); - break; - case 2: - sprintf(vector_type, "%s%d", type, 4); - break; - case 3: - sprintf(vector_type, "%s%d", type, 8); - break; - case 4: - sprintf(vector_type, "%s%d", type, 16); - break; + case 0: sprintf(vector_type, "%s", type); break; + case 1: sprintf(vector_type, "%s%d", type, 2); break; + case 2: sprintf(vector_type, "%s%d", type, 4); break; + case 3: sprintf(vector_type, "%s%d", type, 8); break; + case 4: sprintf(vector_type, "%s%d", type, 16); break; } - if(useWorkgroupReserve == 1){ + if (useWorkgroupReserve == 1) + { createKernelSourceWorkGroup(sourceCode[i], vector_type); } - else if(useSubgroupReserve == 1){ + else if (useSubgroupReserve == 1) + { createKernelSourceSubGroup(sourceCode[i], vector_type); } - else if(useConvenienceBuiltIn == 1){ + else if (useConvenienceBuiltIn == 1) + { createKernelSourceConvenience(sourceCode[i], vector_type); } - else{ + else + { createKernelSource(sourceCode[i], vector_type); } + std::string kernel_source = sourceCode[i].str(); + const char *sources[] = { kernel_source.c_str() }; // Create producer kernel - err = create_single_kernel_helper_with_build_options(context, &program[i], &kernel[ii], 1, (const char**)&sourceCode[i], kernelName[ii], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - align_free( outptr[i] ); - print_error(err, "Error creating program\n"); - return -1; - } - //Create consumer kernel + err = create_single_kernel_helper_with_build_options( + context, &program[i], &kernel[ii], 1, sources, kernelName[ii], + "-cl-std=CL2.0"); + + test_error_ret(err, " Error creating program", -1); + + // Create consumer kernel kernel[ii + 1] = clCreateKernel(program[i], kernelName[ii + 1], &err); - if( kernel[ii + 1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - align_free( outptr[i] ); - log_error("Creating program for %s\n", type); - print_error( err, "Unable to create kernel" ); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); - err = clSetKernelArg(kernel[ii], 0, sizeof(cl_mem), (void*)&buffers[ii]); - err |= clSetKernelArg(kernel[ii], 1, sizeof(cl_mem), (void*)&pipes[i]); - err |= clSetKernelArg(kernel[ii + 1], 0, sizeof(cl_mem), (void*)&pipes[i]); - err |= clSetKernelArg(kernel[ii + 1], 1, sizeof(cl_mem), (void*)&buffers[ii + 1]); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - print_error(err, " clSetKernelArg failed"); - return -1; - } + err = + clSetKernelArg(kernel[ii], 0, sizeof(cl_mem), (void *)&buffers[ii]); + err |= clSetKernelArg(kernel[ii], 1, sizeof(cl_mem), (void *)&pipes[i]); + err |= clSetKernelArg(kernel[ii + 1], 0, sizeof(cl_mem), + (void *)&pipes[i]); + err |= clSetKernelArg(kernel[ii + 1], 1, sizeof(cl_mem), + (void *)&buffers[ii + 1]); + test_error_ret(err, " clSetKernelArg failed", -1); - if(useWorkgroupReserve == 1 || useSubgroupReserve == 1) + if (useWorkgroupReserve == 1 || useSubgroupReserve == 1) { - err = get_max_common_work_group_size( context, kernel[ii], global_work_size[0], &local_work_size[0] ); - test_error( err, "Unable to get work group size to use" ); + err = get_max_common_work_group_size( + context, kernel[ii], global_work_size[0], &local_work_size[0]); + test_error(err, "Unable to get work group size to use"); // Launch Producer kernel - err = clEnqueueNDRangeKernel( queue, kernel[ii], 1, NULL, global_work_size, local_work_size, 0, NULL, &producer_sync_event[i] ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseEvent(producer_sync_event[i]); - clReleaseEvent(consumer_sync_event[i]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - return -1; - } + err = clEnqueueNDRangeKernel(queue, kernel[ii], 1, NULL, + global_work_size, local_work_size, 0, + NULL, &producer_sync_event[i]); + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); } else { // Launch Producer kernel - err = clEnqueueNDRangeKernel( queue, kernel[ii], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event[i] ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseEvent(producer_sync_event[i]); - clReleaseEvent(consumer_sync_event[i]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - return -1; - } + err = clEnqueueNDRangeKernel(queue, kernel[ii], 1, NULL, + global_work_size, NULL, 0, NULL, + &producer_sync_event[i]); + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); } - if(useWorkgroupReserve == 1 || useSubgroupReserve == 1) + if (useWorkgroupReserve == 1 || useSubgroupReserve == 1) { - err = get_max_common_work_group_size( context, kernel[ii + 1], global_work_size[0], &local_work_size[0] ); - test_error( err, "Unable to get work group size to use" ); + err = get_max_common_work_group_size(context, kernel[ii + 1], + global_work_size[0], + &local_work_size[0]); + test_error(err, "Unable to get work group size to use"); // Launch Consumer kernel - err = clEnqueueNDRangeKernel( queue, kernel[ii + 1], 1, NULL, global_work_size, local_work_size, 1, &producer_sync_event[i], &consumer_sync_event[i] ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseEvent(producer_sync_event[i]); - clReleaseEvent(consumer_sync_event[i]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - return -1; - } + err = clEnqueueNDRangeKernel(queue, kernel[ii + 1], 1, NULL, + global_work_size, local_work_size, 1, + &producer_sync_event[i], + &consumer_sync_event[i]); + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); } else { // Launch Consumer kernel - err = clEnqueueNDRangeKernel( queue, kernel[ii + 1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event[i], &consumer_sync_event[i] ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseEvent(producer_sync_event[i]); - clReleaseEvent(consumer_sync_event[i]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - return -1; - } + err = clEnqueueNDRangeKernel( + queue, kernel[ii + 1], 1, NULL, global_work_size, NULL, 1, + &producer_sync_event[i], &consumer_sync_event[i]); + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); } - err = clEnqueueReadBuffer(queue, buffers[ii+1], true, 0, ptrSizes[i]*num_elements, outptr[i], 1, &consumer_sync_event[i], NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[ii]); - clReleaseMemObject(buffers[ii+1]); - clReleaseMemObject(pipes[i]); - clReleaseKernel(kernel[ii]); - clReleaseKernel(kernel[ii+1]); - clReleaseEvent(producer_sync_event[i]); - clReleaseEvent(consumer_sync_event[i]); - clReleaseProgram(program[i]); - align_free(outptr[i]); - return -1; - } + err = clEnqueueReadBuffer(queue, buffers[ii + 1], true, 0, + ptrSizes[i] * num_elements, outptr[i], 1, + &consumer_sync_event[i], NULL); + test_error_ret(err, " clEnqueueReadBuffer failed", -1); - if( fn( inptr[i], outptr[i], (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]))){ - log_error("%s%d test failed\n", type, 1< BufferInPtr; + BufferOwningPtr BufferOutPtr; + clProgramWrapper program; + clKernelWrapper kernel[2]; + size_t size = sizeof(TestStruct); + size_t global_work_size[3]; + cl_int err; + int total_errors = 0; + int i; + MTdataHolder d(gRandomSeed); + clEventWrapper producer_sync_event = NULL; + clEventWrapper consumer_sync_event = NULL; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); global_work_size[0] = (size_t)num_elements; inptr = (TestStruct *)align_malloc(size * num_elements, min_alignment); - for ( i = 0; i < num_elements; i++ ){ + for (i = 0; i < num_elements; i++) + { inptr[i].a = (char)genrand_int32(d); inptr[i].b = genrand_int32(d); } + BufferInPtr.reset(inptr, nullptr, 0, size, true); + buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size * num_elements, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + outptr = align_malloc( size * num_elements, min_alignment); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); + buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size * num_elements, outptr, &err); - if (err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, size, num_elements, NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - clReleaseMemObject(pipe); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); + // Create producer kernel err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, &kernelCode, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - align_free(outptr); - log_error(" Error creating program for struct\n"); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - align_free(outptr); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]); - if (err != CL_SUCCESS){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event ); - if (err != CL_SUCCESS){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); // Launch Consumer kernel err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event ); - if (err != CL_SUCCESS){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size*num_elements, outptr, 1, &consumer_sync_event, NULL); - if (err != CL_SUCCESS){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if( verify_readwrite_struct( inptr, outptr, num_elements)){ log_error("struct_readwrite test failed\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); return -1; } else { log_info("struct_readwrite test passed\n"); } - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); return 0; } diff --git a/test_conformance/pipes/test_pipe_readwrite_errors.cpp b/test_conformance/pipes/test_pipe_readwrite_errors.cpp index cdfaf5e8..1b9fc313 100644 --- a/test_conformance/pipes/test_pipe_readwrite_errors.cpp +++ b/test_conformance/pipes/test_pipe_readwrite_errors.cpp @@ -64,23 +64,26 @@ const char* pipe_readwrite_errors_kernel_code = { int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipe; - cl_mem buffers[3]; - void *outptr; - cl_int *inptr; - cl_program program; - cl_kernel kernel[2]; - size_t global_work_size[3]; - cl_int err; - cl_int size; - cl_int i; - cl_int status = 0; - cl_event producer_sync_event; - cl_event consumer_sync_event; - MTdata d = init_genrand( gRandomSeed ); - const char* kernelName[] = {"test_pipe_write_error", "test_pipe_read_error"}; + clMemWrapper pipe; + clMemWrapper buffers[3]; + void *outptr; + cl_int *inptr; + clProgramWrapper program; + clKernelWrapper kernel[2]; + size_t global_work_size[3]; + cl_int err; + cl_int size; + cl_int i; + cl_int status = 0; + clEventWrapper producer_sync_event; + clEventWrapper consumer_sync_event; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr; + MTdataHolder d(gRandomSeed); + const char *kernelName[] = { "test_pipe_write_error", + "test_pipe_read_error" }; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); global_work_size[0] = num_elements; @@ -88,69 +91,36 @@ int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_com inptr = (cl_int *)align_malloc(size, min_alignment); - for(i = 0; i < (cl_int)(size / sizeof(int)); i++){ + for (i = 0; i < num_elements; i++) + { inptr[i] = (int)genrand_int32(d); } + BufferInPtr.reset(inptr, nullptr, 0, size, true); + + buffers[0] = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); + test_error_ret(err, " clCreateBuffer failed", -1); - buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } outptr = align_malloc(size, min_alignment); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); + buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + buffers[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - align_free( outptr ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); + //Pipe created with max_packets less than global size pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements - (num_elements/2), NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - align_free( outptr ); - clReleaseMemObject(pipe); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); // Create producer kernel err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_readwrite_errors_kernel_code, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - align_free(outptr); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - align_free(outptr); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); @@ -158,49 +128,15 @@ int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_com err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]); err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buffers[2]); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - print_error(err, " clSetKernelArg failed"); - return -1; - } + + test_error_ret(err, " clSetKernelArg failed", -1); // Launch Consumer kernel for empty pipe err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if(status == 0){ log_error("test_pipe_readwrite_errors failed\n"); @@ -212,34 +148,13 @@ int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_com // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &producer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); - if(status == 0){ + if (status == 0) + { log_error("test_pipe_readwrite_errors failed\n"); return -1; } @@ -247,66 +162,27 @@ int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_com status = 0; } + // We will reuse this variable so release the previous referred event. + clReleaseEvent(consumer_sync_event); + // Launch Consumer kernel err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); - if(status == 0) + if (status == 0) { log_error("test_pipe_readwrite_errors failed\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); return -1; } + else + { + status = 0; + } log_info("test_pipe_readwrite_errors passed\n"); - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); + return 0; } diff --git a/test_conformance/pipes/test_pipe_subgroups.cpp b/test_conformance/pipes/test_pipe_subgroups.cpp index 35519b05..b41170ca 100644 --- a/test_conformance/pipes/test_pipe_subgroups.cpp +++ b/test_conformance/pipes/test_pipe_subgroups.cpp @@ -88,30 +88,35 @@ static int verify_result(void *ptr1, void *ptr2, int n) int test_pipe_subgroups_divergence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem pipe; - cl_mem buffers[3]; - cl_int *outptr; - cl_int *inptr; - cl_int *active_work_item_buffer; - cl_program program; - cl_kernel kernel[2]; - size_t global_work_size[3]; - size_t local_work_size[3]; - cl_int err; - cl_int size; - int i; - size_t subgroup_count; - cl_event producer_sync_event = NULL; - cl_event consumer_sync_event = NULL; - const char* kernelName[] = {"test_pipe_subgroups_divergence_write", "test_pipe_subgroups_divergence_read"}; + clMemWrapper pipe; + clMemWrapper buffers[3]; + cl_int *outptr; + cl_int *inptr; + cl_int *active_work_item_buffer; + clProgramWrapper program; + clKernelWrapper kernel[2]; + size_t global_work_size[3]; + size_t local_work_size[3]; + cl_int err; + cl_int size; + int i; + size_t subgroup_count; + clEventWrapper producer_sync_event = NULL; + clEventWrapper consumer_sync_event = NULL; + BufferOwningPtr BufferInPtr; + BufferOwningPtr BufferOutPtr; + const char *kernelName[] = { "test_pipe_subgroups_divergence_write", + "test_pipe_subgroups_divergence_read" }; - size_t min_alignment = get_min_alignment(context); + size_t min_alignment = get_min_alignment(context); global_work_size[0] = (cl_uint)num_elements; - if(!is_extension_available(deviceID, "cl_khr_subgroups")) + if (!is_extension_available(deviceID, "cl_khr_subgroups")) { - log_info("cl_khr_subgroups is not supported on this platoform. Skipping test.\n"); + log_info( + "cl_khr_subgroups is not supported on this platoform. Skipping " + "test.\n"); return CL_SUCCESS; } @@ -125,215 +130,77 @@ int test_pipe_subgroups_divergence(cl_device_id deviceID, cl_context context, cl outptr[i] = 0; active_work_item_buffer[i] = 0; } + BufferInPtr.reset(inptr, nullptr, 0, size, true); + BufferOutPtr.reset(outptr, nullptr, 0, size, true); buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err); - if(err){ - clReleaseMemObject(buffers[0]); - print_error(err, " clCreateBuffer failed\n"); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); buffers[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, outptr, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - align_free( outptr ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); buffers[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, active_work_item_buffer, &err); - if ( err ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - align_free( outptr ); - print_error(err, " clCreateBuffer failed\n" ); - return -1; - } + test_error_ret(err, " clCreateBuffer failed", -1); pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements, NULL, &err); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - align_free( outptr ); - clReleaseMemObject(pipe); - print_error(err, " clCreatePipe failed\n"); - return -1; - } + test_error_ret(err, " clCreatePipe failed", -1); // Create producer kernel err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_subgroups_kernel_code, kernelName[0], "-cl-std=CL2.0"); - if(err){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - align_free(outptr); - print_error(err, "Error creating program\n"); - return -1; - } + test_error_ret(err, " Error creating program", -1); + //Create consumer kernel kernel[1] = clCreateKernel(program, kernelName[1], &err); - if( kernel[1] == NULL || err != CL_SUCCESS) - { - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - align_free(outptr); - print_error(err, "Error creating kernel\n"); - return -1; - } + test_error_ret(err, " Error creating kernel", -1); err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]); err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buffers[2]); err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe); err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]); - if ( err != CL_SUCCESS ){ - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - print_error(err, " clSetKernelArg failed"); - return -1; - } + test_error_ret(err, " clSetKernelArg failed", -1); err = get_max_common_work_group_size( context, kernel[0], global_work_size[0], &local_work_size[0] ); - if( err != CL_SUCCESS) - { - test_error( err, "Unable to get work group size to use" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " Unable to get work group size to use", -1); - cl_platform_id platform; - err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform), &platform, NULL); - clGetKernelSubGroupInfoKHR_fn clGetKernelSubGroupInfoKHR = (clGetKernelSubGroupInfoKHR_fn) clGetExtensionFunctionAddressForPlatform(platform, "clGetKernelSubGroupInfoKHR"); + cl_platform_id platform; + err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform), + &platform, NULL); + test_error_ret(err, " clGetDeviceInfo failed", -1); + + clGetKernelSubGroupInfoKHR_fn clGetKernelSubGroupInfoKHR = + (clGetKernelSubGroupInfoKHR_fn)clGetExtensionFunctionAddressForPlatform( + platform, "clGetKernelSubGroupInfoKHR"); err = clGetKernelSubGroupInfoKHR(kernel[0], deviceID, CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR, sizeof(local_work_size[0]), &local_work_size[0], sizeof(subgroup_count), &subgroup_count, NULL); + test_error_ret(err, " clGetKernelSubGroupInfoKHR failed", -1); if(subgroup_count <= 1) { log_info("Only 1 subgroup per workgroup for the kernel. Hence no divergence among subgroups possible. Skipping test.\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program); - align_free(outptr); return CL_SUCCESS; } // Launch Producer kernel err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, local_work_size, 0, NULL, &producer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[2], true, 0, size, active_work_item_buffer, 1, &producer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } - + test_error_ret(err, " clEnqueueReadBuffer failed", -1); // Launch Consumer kernel err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, local_work_size, 1, &producer_sync_event, &consumer_sync_event ); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueNDRangeKernel failed", -1); err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL); - if ( err != CL_SUCCESS ){ - print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); - return -1; - } + test_error_ret(err, " clEnqueueReadBuffer failed", -1); if( verify_result( active_work_item_buffer, outptr, num_elements)){ log_error("test_pipe_subgroups_divergence failed\n"); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); return -1; } else { log_info("test_pipe_subgroups_divergence passed\n"); } - //cleanup - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseMemObject(buffers[2]); - clReleaseMemObject(pipe); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseEvent(producer_sync_event); - clReleaseEvent(consumer_sync_event); - clReleaseProgram(program); - align_free(outptr); return 0; }