From cd99c874b23264e5187d5ef5eb442cb8e3cfe7b2 Mon Sep 17 00:00:00 2001 From: ellnor01 <51320439+ellnor01@users.noreply.github.com> Date: Fri, 18 Dec 2020 07:50:30 +0000 Subject: [PATCH] Reduce number of compilations in buffer suite (#1082) * Reduce number of compilations in buffer suite Extracts program and kernel compilation from mem_flags loop as they were being recompiled unnecessarily. Fixes #1020 Signed-off-by: Ellen Norris-Thompson * Remove misplaced frees in buffer tests Contributes #1020 Signed-off-by: Ellen Norris-Thompson --- test_conformance/buffers/test_buffer_fill.cpp | 213 +++++++----------- test_conformance/buffers/test_buffer_map.cpp | 50 ++-- test_conformance/buffers/test_buffer_read.cpp | 167 ++++++-------- .../buffers/test_buffer_write.cpp | 137 +++++------ 4 files changed, 215 insertions(+), 352 deletions(-) diff --git a/test_conformance/buffers/test_buffer_fill.cpp b/test_conformance/buffers/test_buffer_fill.cpp index 5c1dd48e..2a12bd8c 100644 --- a/test_conformance/buffers/test_buffer_fill.cpp +++ b/test_conformance/buffers/test_buffer_fill.cpp @@ -562,11 +562,11 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu int loops, void *inptr[5], void *hostptr[5], void *pattern[5], size_t offset_elements, size_t fill_elements, const char *kernelCode[], const char *kernelName[], int (*fn)(void *,void *,int) ) { - cl_mem buffers[10]; + clMemWrapper buffers[10]; void *outptr[5]; - cl_program program[5]; - cl_kernel kernel[5]; - cl_event event[2]; + clProgramWrapper program[5]; + clKernelWrapper kernel[5]; + clEventWrapper event[2]; size_t ptrSizes[5]; size_t global_work_size[3]; int err; @@ -584,12 +584,22 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu ptrSizes[3] = ptrSizes[2] << 1; ptrSizes[4] = ptrSizes[3] << 1; - for (src_flag_id=0; src_flag_id < NUM_FLAGS; src_flag_id++) { - log_info("Testing with cl_mem_flags: %s\n", flag_set_names[src_flag_id]); + loops = (loops < 5 ? loops : 5); + for (i = 0; i < loops; i++) + { + ii = i << 1; + + err = create_single_kernel_helper(context, &program[i], &kernel[i], 1, + &kernelCode[i], kernelName[i]); + if (err) + { + log_error(" Error creating program for %s\n", type); + return -1; + } + + for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) + { - loops = ( loops < 5 ? loops : 5 ); - for ( i = 0; i < loops; i++ ){ - ii = i << 1; if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) buffers[ii] = clCreateBuffer(context, flag_set[src_flag_id], ptrSizes[i] * num_elements, hostptr[i], &err); else @@ -612,7 +622,6 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu buffers[ii+1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, ptrSizes[i] * num_elements, outptr[i], &err); if ( !buffers[ii+1] || err){ print_error(err, "clCreateBuffer failed\n" ); - clReleaseMemObject( buffers[ii] ); align_free( outptr[i] ); return -1; } @@ -625,17 +634,6 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu */ if ( err != CL_SUCCESS ){ print_error( err, " clEnqueueFillBuffer failed" ); - clReleaseMemObject( buffers[ii] ); - clReleaseMemObject( buffers[ii+1] ); - align_free( outptr[i] ); - return -1; - } - - err = create_single_kernel_helper( context, &program[i], &kernel[i], 1, &kernelCode[i], kernelName[i] ); - if ( err ){ - log_error( " Error creating program for %s\n", type ); - clReleaseMemObject( buffers[ii] ); - clReleaseMemObject( buffers[ii+1] ); align_free( outptr[i] ); return -1; } @@ -644,10 +642,6 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu err |= clSetKernelArg( kernel[i], 1, sizeof( cl_mem ), (void *)&buffers[ii+1] ); if ( err != CL_SUCCESS ){ print_error( err, "clSetKernelArg failed" ); - clReleaseKernel( kernel[i] ); - clReleaseProgram( program[i] ); - clReleaseMemObject( buffers[ii] ); - clReleaseMemObject( buffers[ii+1] ); align_free( outptr[i] ); return -1; } @@ -655,14 +649,9 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu err = clWaitForEvents( 1, &(event[0]) ); if ( err != CL_SUCCESS ){ print_error( err, "clWaitForEvents() failed" ); - clReleaseKernel( kernel[i] ); - clReleaseProgram( program[i] ); - clReleaseMemObject( buffers[ii] ); - clReleaseMemObject( buffers[ii+1] ); align_free( outptr[i] ); return -1; } - clReleaseEvent(event[0]); err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, global_work_size, NULL, 0, NULL, NULL ); if (err != CL_SUCCESS){ @@ -680,21 +669,18 @@ int test_buffer_fill( cl_device_id deviceID, cl_context context, cl_command_queu if ( err != CL_SUCCESS ){ print_error( err, "clWaitForEvents() failed" ); } - clReleaseEvent(event[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<a = (cl_int)genrand_int32(d); - pattern->b = (cl_float)get_random_float( -FLT_MAX, FLT_MAX, d ); + for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) + { + log_info("Testing with cl_mem_flags: %s\n", + flag_set_names[src_flag_id]); - inptr = (TestStruct *)align_malloc(ptrSize * num_elements, min_alignment); - for ( j = 0; j < offset_elements; j++ ) { - inptr[j].a = 0; - inptr[j].b =0; - } - for ( j = offset_elements; j < offset_elements + fill_elements; j++ ) { - inptr[j].a = pattern->a; - inptr[j].b = pattern->b; - } - for ( j = offset_elements + fill_elements; j < (size_t)num_elements; j++ ) { - inptr[j].a = 0; - inptr[j].b = 0; + err = create_single_kernel_helper(context, &program, &kernel, 1, + &struct_kernel_code, + "read_fill_struct"); + if (err) + { + log_error(" Error creating program for struct\n"); + free_mtdata(d); + return -1; } - hostptr = (TestStruct *)align_malloc(ptrSize * num_elements, min_alignment); - memset(hostptr, 0, ptrSize * num_elements); + // Test with random offsets and fill sizes + for (n = 0; n < 8; n++) + { + offset_elements = + (size_t)get_random_float(0.f, (float)(num_elements - 8), d); + fill_elements = (size_t)get_random_float( + 8.f, (float)(num_elements - offset_elements), d); + log_info("Testing random fill from offset %d for %d elements: \n", + (int)offset_elements, (int)fill_elements); - for (src_flag_id=0; src_flag_id < NUM_FLAGS; src_flag_id++) { - log_info("Testing with cl_mem_flags: %s\n", flag_set_names[src_flag_id]); + pattern.a = (cl_int)genrand_int32(d); + pattern.b = (cl_float)get_random_float(-FLT_MAX, FLT_MAX, d); + + inptr = (TestStruct *)align_malloc(ptrSize * num_elements, + min_alignment); + for (j = 0; j < offset_elements; j++) + { + inptr[j].a = 0; + inptr[j].b = 0; + } + for (j = offset_elements; j < offset_elements + fill_elements; j++) + { + inptr[j].a = pattern.a; + inptr[j].b = pattern.b; + } + for (j = offset_elements + fill_elements; j < (size_t)num_elements; + j++) + { + inptr[j].a = 0; + inptr[j].b = 0; + } + + hostptr = (TestStruct *)align_malloc(ptrSize * num_elements, + min_alignment); + memset(hostptr, 0, ptrSize * num_elements); if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) buffers[0] = clCreateBuffer(context, flag_set[src_flag_id], ptrSize * num_elements, hostptr, &err); @@ -762,9 +770,6 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma buffers[0] = clCreateBuffer(context, flag_set[src_flag_id], ptrSize * num_elements, NULL, &err); if ( err ){ print_error(err, " clCreateBuffer failed\n" ); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -774,9 +779,6 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma err = clEnqueueWriteBuffer(queue, buffers[0], CL_FALSE, 0, ptrSize * num_elements, hostptr, 0, NULL, NULL); if ( err != CL_SUCCESS ){ print_error(err, " clEnqueueWriteBuffer failed\n" ); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -789,45 +791,21 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma if ( ! buffers[1] || err){ print_error(err, " clCreateBuffer failed\n" ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); return -1; } - err = clEnqueueFillBuffer(queue, buffers[0], pattern, ptrSize, - ptrSize * offset_elements, ptrSize * fill_elements, - 0, NULL, &(event[0])); + err = clEnqueueFillBuffer( + queue, buffers[0], &pattern, ptrSize, ptrSize * offset_elements, + ptrSize * fill_elements, 0, NULL, &(event[0])); /* uncomment for test debugging err = clEnqueueWriteBuffer(queue, buffers[0], CL_FALSE, 0, ptrSize * num_elements, inptr, 0, NULL, &(event[0])); */ if ( err != CL_SUCCESS ){ print_error( err, " clEnqueueFillBuffer failed" ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); - align_free( (void *)inptr ); - align_free( (void *)hostptr ); - free_mtdata(d); - return -1; - } - - err = create_single_kernel_helper( context, &program, &kernel, 1, &struct_kernel_code, "read_fill_struct" ); - if ( err ){ - log_error( " Error creating program for struct\n" ); - align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -838,14 +816,7 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma err |= clSetKernelArg( kernel, 1, sizeof( cl_mem ), (void *)&buffers[1] ); if ( err != CL_SUCCESS ){ print_error( err, " clSetKernelArg failed" ); - clReleaseKernel( kernel ); - clReleaseProgram( program ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -855,14 +826,7 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma err = clWaitForEvents( 1, &(event[0]) ); if ( err != CL_SUCCESS ){ print_error( err, "clWaitForEvents() failed" ); - clReleaseKernel( kernel ); - clReleaseProgram( program ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -873,14 +837,7 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL ); if ( err != CL_SUCCESS ){ print_error( err, " clEnqueueNDRangeKernel failed" ); - clReleaseKernel( kernel ); - clReleaseProgram( program ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -890,14 +847,7 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma err = clEnqueueReadBuffer( queue, buffers[1], CL_FALSE, 0, ptrSize * num_elements, outptr, 0, NULL, &(event[1]) ); if ( err != CL_SUCCESS ){ print_error( err, " clEnqueueReadBuffer failed" ); - clReleaseKernel( kernel ); - clReleaseProgram( program ); align_free( outptr ); - clReleaseMemObject(buffers[0]); - clReleaseMemObject(buffers[1]); - clReleaseEvent( event[0] ); - clReleaseEvent( event[1] ); - free( (void *)pattern ); align_free( (void *)inptr ); align_free( (void *)hostptr ); free_mtdata(d); @@ -918,15 +868,10 @@ int test_buffer_fill_struct( cl_device_id deviceID, cl_context context, cl_comma log_info( " buffer_FILL async struct test passed\n" ); } // cleanup - clReleaseKernel( kernel ); - clReleaseProgram( program ); align_free( outptr ); - clReleaseMemObject( buffers[0] ); - clReleaseMemObject( buffers[1] ); + align_free((void *)inptr); + align_free((void *)hostptr); } // src cl_mem_flag - free( (void *)pattern ); - align_free( (void *)inptr ); - align_free( (void *)hostptr ); } free_mtdata(d); diff --git a/test_conformance/buffers/test_buffer_map.cpp b/test_conformance/buffers/test_buffer_map.cpp index f0363dd5..3cbcd387 100644 --- a/test_conformance/buffers/test_buffer_map.cpp +++ b/test_conformance/buffers/test_buffer_map.cpp @@ -554,10 +554,10 @@ static int verify_read_struct( void *ptr, int n ) static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops, const char *kernelCode[], const char *kernelName[], int (*fn)(void *,int) ) { - cl_mem buffers[5]; + clMemWrapper buffers[5]; void *outptr[5]; - cl_program program[5]; - cl_kernel kernel[5]; + clProgramWrapper program[5]; + clKernelWrapper kernel[5]; size_t threads[3], localThreads[3]; cl_int err; int i; @@ -580,10 +580,20 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c if (! gHasLong && strstr(type,"long")) return 0; - for (src_flag_id=0; src_flag_id < NUM_FLAGS; src_flag_id++) { - log_info("Testing with cl_mem_flags src: %s\n", flag_set_names[src_flag_id]); + for (i = 0; i < loops; i++) + { + + err = create_single_kernel_helper(context, &program[i], &kernel[i], 1, + &kernelCode[i], kernelName[i]); + if (err) + { + log_error(" Error creating program for %s\n", type); + return -1; + } + + for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) + { - for ( i = 0; i < loops; i++ ){ outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment); if ( ! outptr[i] ){ log_error( " unable to allocate %d bytes of memory\n", (int)ptrSizes[i] * num_elements ); @@ -602,20 +612,9 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c return -1; } - err = create_single_kernel_helper(context, &program[i], &kernel[i], 1, &kernelCode[i], kernelName[i] ); - if ( err ){ - log_error( " Error creating program for %s\n", type ); - clReleaseMemObject( buffers[i] ); - align_free( outptr[i] ); - return -1; - } - err = clSetKernelArg( kernel[i], 0, sizeof( cl_mem ), (void *)&buffers[i] ); if ( err != CL_SUCCESS ){ print_error( err, "clSetKernelArg failed\n" ); - clReleaseKernel( kernel[i] ); - clReleaseProgram( program[i] ); - clReleaseMemObject( buffers[i] ); align_free( outptr[i] ); return -1; } @@ -628,9 +627,6 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, localThreads, 0, NULL, NULL ); if ( err != CL_SUCCESS ){ print_error( err, "clEnqueueNDRangeKernel failed\n" ); - clReleaseKernel( kernel[i] ); - clReleaseProgram( program[i] ); - clReleaseMemObject( buffers[i] ); align_free( outptr[i] ); return -1; } @@ -638,29 +634,23 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c mappedPtr = clEnqueueMapBuffer(queue, buffers[i], CL_TRUE, CL_MAP_READ, 0, ptrSizes[i]*num_elements, 0, NULL, NULL, &err); if ( err != CL_SUCCESS ){ print_error( err, "clEnqueueMapBuffer failed" ); - clReleaseKernel( kernel[i] ); - clReleaseProgram( program[i] ); - clReleaseMemObject( buffers[i] ); align_free( outptr[i] ); return -1; } if (fn(mappedPtr, num_elements*(1<