From 25d9ff5d6ec4dd8140b74f093b888984aa7580ee Mon Sep 17 00:00:00 2001 From: ellnor01 <51320439+ellnor01@users.noreply.github.com> Date: Thu, 7 Jan 2021 11:34:42 +0000 Subject: [PATCH] Using helper functions for clCreateKernel (#1064) * Using helper functions for clCreateKernel Uses of clCreateKernel following create program helper functions, have been incorporated into create_single_kernel_helper when suitable. Contributes #31 Signed-off-by: Ellen Norris-Thompson * Skip tests using clCompileProgram in offline mode Contributes #31 Signed-off-by: Ellen Norris-Thompson * Using type wrappers when using kernel helper functions Also includes fix for windows build Fixes #31 Signed-off-by: Ellen Norris-Thompson * Remove clReleaseKernel for wrapped kernel Fixes #31 Signed-off-by: Ellen Norris-Thompson --- test_common/harness/errorHelpers.cpp | 2 + test_common/harness/imageHelpers.cpp | 42 +----- test_conformance/api/test_create_kernels.cpp | 9 +- test_conformance/api/test_null_buffer_arg.cpp | 11 +- test_conformance/api/test_retain.cpp | 8 +- test_conformance/api/test_retain_program.cpp | 9 +- test_conformance/d3d10/harness.cpp | 33 +--- test_conformance/d3d11/harness.cpp | 37 +---- test_conformance/math_brute_force/main.cpp | 142 ++++++++---------- test_conformance/printf/test_printf.cpp | 19 +-- ...l_khr_spirv_no_integer_wrap_decoration.cpp | 10 +- test_conformance/spirv_new/test_op_fmath.cpp | 10 +- .../spirv_new/test_op_vector_times_scalar.cpp | 10 +- 13 files changed, 111 insertions(+), 231 deletions(-) diff --git a/test_common/harness/errorHelpers.cpp b/test_common/harness/errorHelpers.cpp index c1d06028..8f3c1883 100644 --- a/test_common/harness/errorHelpers.cpp +++ b/test_common/harness/errorHelpers.cpp @@ -686,6 +686,8 @@ const char *subtests_to_skip_with_offline_compiler[] = { "unload_build_info", "unload_program_binaries", "features_macro", + "progvar_prog_scope_misc", + "library_function" }; int check_functions_for_offline_compiler(const char *subtestname, diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp index 26110a47..b785f64d 100644 --- a/test_common/harness/imageHelpers.cpp +++ b/test_common/harness/imageHelpers.cpp @@ -2917,7 +2917,7 @@ int DetectFloatToHalfRoundingMode( } // Create our program, and a kernel - const char *kernel[1] = { + const char *kernelSource[1] = { "kernel void detect_round( global float4 *in, write_only image2d_t " "out )\n" "{\n" @@ -2927,8 +2927,9 @@ int DetectFloatToHalfRoundingMode( }; clProgramWrapper program; - err = create_single_kernel_helper_create_program(context, &program, 1, - kernel); + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &program, &kernel, 1, + kernelSource, "detect_round"); if (NULL == program || err) { @@ -2953,29 +2954,7 @@ int DetectFloatToHalfRoundingMode( return err; } - err = clBuildProgram(program, 1, &device, "", NULL, NULL); - if (err) - { - log_error("Error: could not build program in " - "DetectFloatToHalfRoundingMode (%d)", - err); - clReleaseMemObject(inBuf); - clReleaseMemObject(outImage); - return err; - } - - cl_kernel k = clCreateKernel(program, "detect_round", &err); - if (NULL == k || err) - { - log_error("Error: could not create kernel in " - "DetectFloatToHalfRoundingMode (%d)", - err); - clReleaseMemObject(inBuf); - clReleaseMemObject(outImage); - return err; - } - - err = clSetKernelArg(k, 0, sizeof(cl_mem), &inBuf); + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &inBuf); if (err) { log_error("Error: could not set argument 0 of kernel in " @@ -2983,11 +2962,10 @@ int DetectFloatToHalfRoundingMode( err); clReleaseMemObject(inBuf); clReleaseMemObject(outImage); - clReleaseKernel(k); return err; } - err = clSetKernelArg(k, 1, sizeof(cl_mem), &outImage); + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &outImage); if (err) { log_error("Error: could not set argument 1 of kernel in " @@ -2995,14 +2973,13 @@ int DetectFloatToHalfRoundingMode( err); clReleaseMemObject(inBuf); clReleaseMemObject(outImage); - clReleaseKernel(k); return err; } // Run the kernel size_t global_work_size = count; - err = clEnqueueNDRangeKernel(q, k, 1, NULL, &global_work_size, NULL, 0, - NULL, NULL); + err = clEnqueueNDRangeKernel(q, kernel, 1, NULL, &global_work_size, + NULL, 0, NULL, NULL); if (err) { log_error("Error: could not enqueue kernel in " @@ -3010,7 +2987,6 @@ int DetectFloatToHalfRoundingMode( err); clReleaseMemObject(inBuf); clReleaseMemObject(outImage); - clReleaseKernel(k); return err; } @@ -3028,7 +3004,6 @@ int DetectFloatToHalfRoundingMode( err); clReleaseMemObject(inBuf); clReleaseMemObject(outImage); - clReleaseKernel(k); return err; } @@ -3083,7 +3058,6 @@ int DetectFloatToHalfRoundingMode( // clean up clReleaseMemObject(inBuf); clReleaseMemObject(outImage); - clReleaseKernel(k); return err; } diff --git a/test_conformance/api/test_create_kernels.cpp b/test_conformance/api/test_create_kernels.cpp index 79e01fdb..568e84cb 100644 --- a/test_conformance/api/test_create_kernels.cpp +++ b/test_conformance/api/test_create_kernels.cpp @@ -525,11 +525,10 @@ int test_repeated_setup_cleanup(cl_device_id deviceID, cl_context context, cl_co local_queue = clCreateCommandQueue(local_context, deviceID, 0, &error); test_error( error, "clCreateCommandQueue failed"); - error = create_single_kernel_helper(local_context, &local_program, NULL, 1, &repeate_test_kernel, NULL); - test_error( error, "Unable to build test program" ); - - local_kernel = clCreateKernel(local_program, "test_kernel", &error); - test_error( error, "clCreateKernel failed"); + error = create_single_kernel_helper( + local_context, &local_program, &local_kernel, 1, + &repeate_test_kernel, "test_kernel"); + test_error(error, "Unable to create kernel"); local_mem_in = clCreateBuffer(local_context, CL_MEM_READ_ONLY, TEST_SIZE*sizeof(cl_int), NULL, &error); test_error( error, "clCreateBuffer failed"); diff --git a/test_conformance/api/test_null_buffer_arg.cpp b/test_conformance/api/test_null_buffer_arg.cpp index ba43f183..d412d4ea 100644 --- a/test_conformance/api/test_null_buffer_arg.cpp +++ b/test_conformance/api/test_null_buffer_arg.cpp @@ -157,14 +157,13 @@ int test_null_buffer_arg(cl_device_id device, cl_context context, // prep kernel: if (gIsEmbedded) - status = create_single_kernel_helper(context, &program, NULL, 1, &kernel_string, NULL); + status = create_single_kernel_helper(context, &program, &kernel, 1, + &kernel_string, "test_kernel"); else - status = create_single_kernel_helper(context, &program, NULL, 1, &kernel_string_long, NULL); + status = create_single_kernel_helper( + context, &program, &kernel, 1, &kernel_string_long, "test_kernel"); - test_error(status, "Unable to build test program"); - - kernel = clCreateKernel(program, "test_kernel", &status); - test_error(status, "CreateKernel failed."); + test_error(status, "Unable to create kernel"); cl_mem dev_src = clCreateBuffer(context, CL_MEM_READ_ONLY, NITEMS*sizeof(cl_float), NULL, NULL); diff --git a/test_conformance/api/test_retain.cpp b/test_conformance/api/test_retain.cpp index cf065bcd..6e66c7da 100644 --- a/test_conformance/api/test_retain.cpp +++ b/test_conformance/api/test_retain.cpp @@ -251,11 +251,9 @@ int test_retain_mem_object_set_kernel_arg(cl_device_id deviceID, cl_context cont err = clSetMemObjectDestructorCallback( buffer, callback, nullptr ); test_error( err, "Unable to set destructor callback" ); - err = create_single_kernel_helper( context, &program, nullptr, 1, testProgram, nullptr ); - test_error( err, "Unable to build sample program" ); - - kernel = clCreateKernel( program, "sample_test", &err ); - test_error( err, "Unable to create sample_test kernel" ); + err = create_single_kernel_helper(context, &program, &kernel, 1, + testProgram, "sample_test"); + test_error(err, "Unable to build sample program and sample_test kernel"); err = clSetKernelArg( kernel, 0, sizeof(cl_mem), &buffer ); test_error( err, "Unable to set kernel argument" ); diff --git a/test_conformance/api/test_retain_program.cpp b/test_conformance/api/test_retain_program.cpp index aa9c8b36..b9fc8b7e 100644 --- a/test_conformance/api/test_retain_program.cpp +++ b/test_conformance/api/test_retain_program.cpp @@ -28,14 +28,11 @@ int test_release_kernel_order(cl_device_id deviceID, cl_context context, cl_comm int error; const char *testProgram[] = { "__kernel void sample_test(__global int *data){}" }; - /* Create a test program */ - error = create_single_kernel_helper(context, &program, NULL, 1, testProgram, NULL); + /* Create a test program and kernel from it */ + error = create_single_kernel_helper(context, &program, &kernel, 1, + testProgram, "sample_test"); test_error( error, "Unable to build sample program to test with" ); - /* And create a kernel from it */ - kernel = clCreateKernel( program, "sample_test", &error ); - test_error( error, "Unable to create kernel" ); - /* Now try freeing the program first, then the kernel. If refcounts are right, this should work just fine */ clReleaseProgram( program ); clReleaseKernel( kernel ); diff --git a/test_conformance/d3d10/harness.cpp b/test_conformance/d3d10/harness.cpp index ffdfea5a..93f2281d 100644 --- a/test_conformance/d3d10/harness.cpp +++ b/test_conformance/d3d10/harness.cpp @@ -367,41 +367,12 @@ cl_int HarnessD3D10_CreateKernelFromSource( const char *sourceTexts[] = {source}; size_t sourceLengths[] = {strlen(source) }; - status = create_single_kernel_helper_create_program(context, &program, 1, &sourceTexts[0]); + status = create_single_kernel_helper(context, &program, &kernel, 1, + &sourceTexts[0], entrypoint); TestRequire( CL_SUCCESS == status, "clCreateProgramWithSource failed"); } - status = clBuildProgram( - program, - 0, - NULL, - NULL, - NULL, - NULL); - if (CL_SUCCESS != status) - { - char log[2048] = {0}; - status = clGetProgramBuildInfo( - program, - device, - CL_PROGRAM_BUILD_LOG, - sizeof(log), - log, - NULL); - TestPrint("error: %s\n", log); - TestRequire( - CL_SUCCESS == status, - "Compilation error log:\n%s\n", log); - } - - kernel = clCreateKernel( - program, - entrypoint, - &status); - TestRequire( - CL_SUCCESS == status, - "clCreateKernel failed"); clReleaseProgram(program); *outKernel = kernel; diff --git a/test_conformance/d3d11/harness.cpp b/test_conformance/d3d11/harness.cpp index 687c6da2..90ba200b 100644 --- a/test_conformance/d3d11/harness.cpp +++ b/test_conformance/d3d11/harness.cpp @@ -400,41 +400,10 @@ cl_int HarnessD3D11_CreateKernelFromSource( const char *sourceTexts[] = {source}; size_t sourceLengths[] = {strlen(source) }; - status = create_single_kernel_helper_create_program(context, &program, 1, &sourceTexts[0]); - TestRequire( - CL_SUCCESS == status, - "clCreateProgramWithSource failed"); + status = create_single_kernel_helper(context, &program, &kernel, 1, + &sourceTexts[0], entrypoint); + TestRequire(CL_SUCCESS == status, "Kernel creation failed"); } - status = clBuildProgram( - program, - 0, - NULL, - NULL, - NULL, - NULL); - if (CL_SUCCESS != status) - { - char log[2048] = {0}; - status = clGetProgramBuildInfo( - program, - device, - CL_PROGRAM_BUILD_LOG, - sizeof(log), - log, - NULL); - TestPrint("error: %s\n", log); - TestRequire( - CL_SUCCESS == status, - "Compilation error log:\n%s\n", log); - } - - kernel = clCreateKernel( - program, - entrypoint, - &status); - TestRequire( - CL_SUCCESS == status, - "clCreateKernel failed"); clReleaseProgram(program); *outKernel = kernel; diff --git a/test_conformance/math_brute_force/main.cpp b/test_conformance/math_brute_force/main.cpp index 8f2e0a0c..d7f2ebf6 100644 --- a/test_conformance/math_brute_force/main.cpp +++ b/test_conformance/math_brute_force/main.cpp @@ -25,6 +25,7 @@ #include "harness/errorHelpers.h" #include "harness/kernelHelpers.h" #include "harness/parseParameters.h" +#include "harness/typeWrappers.h" #if defined( __APPLE__ ) #include @@ -1384,36 +1385,36 @@ void _LogBuildError( cl_program p, int line, const char *file ) int InitILogbConstants( void ) { int error; - const char *kernel = - "__kernel void GetILogBConstants( __global int *out )\n" - "{\n" - " out[0] = FP_ILOGB0;\n" - " out[1] = FP_ILOGBNAN;\n" - "}\n"; + const char *kernelSource = + R"(__kernel void GetILogBConstants( __global int *out ) + { + out[0] = FP_ILOGB0; + out[1] = FP_ILOGBNAN; + })"; - cl_program query; - error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL); - if (NULL == query || error) + clProgramWrapper query; + clKernelWrapper kernel; + error = create_single_kernel_helper(gContext, &query, &kernel, 1, + &kernelSource, "GetILogBConstants"); + if (error != CL_SUCCESS) { - vlog_error( "Error: Unable to create program to get FP_ILOGB0 and FP_ILOGBNAN for the device. (%d)", error ); + vlog_error("Error: Unable to create kernel to get FP_ILOGB0 and " + "FP_ILOGBNAN for the device. (%d)", + error); return error; } - cl_kernel k = clCreateKernel( query, "GetILogBConstants", &error ); - if( NULL == k || error) - { - vlog_error( "Error: Unable to create kernel to get FP_ILOGB0 and FP_ILOGBNAN for the device. Err = %d", error ); - return error; - } - - if((error = clSetKernelArg(k, 0, sizeof( gOutBuffer[gMinVectorSizeIndex]), &gOutBuffer[gMinVectorSizeIndex]))) + if ((error = + clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]), + &gOutBuffer[gMinVectorSizeIndex]))) { vlog_error( "Error: Unable to set kernel arg to get FP_ILOGB0 and FP_ILOGBNAN for the device. Err = %d", error ); return error; } size_t dim = 1; - if((error = clEnqueueNDRangeKernel(gQueue, k, 1, NULL, &dim, NULL, 0, NULL, NULL) )) + if ((error = clEnqueueNDRangeKernel(gQueue, kernel, 1, NULL, &dim, NULL, 0, + NULL, NULL))) { vlog_error( "Error: Unable to execute kernel to get FP_ILOGB0 and FP_ILOGBNAN for the device. Err = %d", error ); return error; @@ -1429,45 +1430,43 @@ int InitILogbConstants( void ) gDeviceILogb0 = data.ilogb0; gDeviceILogbNaN = data.ilogbnan; - clReleaseKernel(k); - clReleaseProgram(query); - return 0; } int IsTininessDetectedBeforeRounding( void ) { int error; - const char *kernel = - "__kernel void IsTininessDetectedBeforeRounding( __global float *out )\n" - "{\n" - " volatile float a = 0x1.000002p-126f;\n" - " volatile float b = 0x1.fffffcp-1f;\n" // product is 0x1.fffffffffff8p-127 - " out[0] = a * b;\n" - "}\n"; + const char *kernelSource = + R"(__kernel void IsTininessDetectedBeforeRounding( __global float *out ) + { + volatile float a = 0x1.000002p-126f; + volatile float b = 0x1.fffffcp-1f; + out[0] = a * b; // product is 0x1.fffffffffff8p-127 + })"; - cl_program query; - error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL); + clProgramWrapper query; + clKernelWrapper kernel; + error = + create_single_kernel_helper(gContext, &query, &kernel, 1, &kernelSource, + "IsTininessDetectedBeforeRounding"); if (error != CL_SUCCESS) { - vlog_error( "Error: Unable to create program to detect how tininess is detected for the device. (%d)", error ); + vlog_error("Error: Unable to create kernel to detect how tininess is " + "detected for the device. (%d)", + error); return error; } - cl_kernel k = clCreateKernel( query, "IsTininessDetectedBeforeRounding", &error ); - if( NULL == k || error) - { - vlog_error( "Error: Unable to create kernel to detect how tininess is detected for the device. Err = %d", error ); - return error; - } - - if((error = clSetKernelArg(k, 0, sizeof( gOutBuffer[gMinVectorSizeIndex]), &gOutBuffer[gMinVectorSizeIndex]))) + if ((error = + clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]), + &gOutBuffer[gMinVectorSizeIndex]))) { vlog_error( "Error: Unable to set kernel arg to detect how tininess is detected for the device. Err = %d", error ); return error; } size_t dim = 1; - if((error = clEnqueueNDRangeKernel(gQueue, k, 1, NULL, &dim, NULL, 0, NULL, NULL) )) + if ((error = clEnqueueNDRangeKernel(gQueue, kernel, 1, NULL, &dim, NULL, 0, + NULL, NULL))) { vlog_error( "Error: Unable to execute kernel to detect how tininess is detected for the device. Err = %d", error ); return error; @@ -1482,9 +1481,6 @@ int IsTininessDetectedBeforeRounding( void ) gCheckTininessBeforeRounding = 0 == (data.f & 0x7fffffff); - clReleaseKernel(k); - clReleaseProgram(query); - return 0; } @@ -1505,22 +1501,11 @@ int MakeKernel(const char **c, cl_uint count, const char *name, cl_kernel *k, strcat(options, " -cl-fast-relaxed-math"); } - error = create_single_kernel_helper(gContext, p, NULL, count, c, NULL, options); + error = + create_single_kernel_helper(gContext, p, k, count, c, name, options); if (error != CL_SUCCESS) { - vlog_error("\t\tFAILED -- Failed to create program. (%d)\n", error); - return error; - } - - *k = clCreateKernel( *p, name, &error ); - if( NULL == *k || error ) - { - char buffer[2048] = ""; - - vlog_error("\t\tFAILED -- clCreateKernel() failed: (%d)\n", error); - clGetProgramBuildInfo(*p, gDevice, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL); - vlog_error("Log: %s\n", buffer); - clReleaseProgram( *p ); + vlog_error("\t\tFAILED -- Failed to create kernel. (%d)\n", error); return error; } @@ -1581,36 +1566,36 @@ int MakeKernels(const char **c, cl_uint count, const char *name, static int IsInRTZMode( void ) { int error; - const char *kernel = - "__kernel void GetRoundingMode( __global int *out )\n" - "{\n" - " volatile float a = 0x1.0p23f;\n" - " volatile float b = -0x1.0p23f;\n" - " out[0] = (a + 0x1.fffffep-1f == a) && (b - 0x1.fffffep-1f == b);\n" - "}\n"; + const char *kernelSource = + R"(__kernel void GetRoundingMode( __global int *out ) + { + volatile float a = 0x1.0p23f; + volatile float b = -0x1.0p23f; + out[0] = (a + 0x1.fffffep-1f == a) && (b - 0x1.fffffep-1f == b); + "})"; - cl_program query; - error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL); + clProgramWrapper query; + clKernelWrapper kernel; + error = create_single_kernel_helper(gContext, &query, &kernel, 1, + &kernelSource, "GetRoundingMode"); if (error != CL_SUCCESS) { - vlog_error( "Error: Unable to create program to detect RTZ mode for the device. (%d)", error ); + vlog_error("Error: Unable to create kernel to detect RTZ mode for the " + "device. (%d)", + error); return error; } - cl_kernel k = clCreateKernel( query, "GetRoundingMode", &error ); - if( NULL == k || error) - { - vlog_error( "Error: Unable to create kernel to gdetect RTZ mode for the device. Err = %d", error ); - return error; - } - - if((error = clSetKernelArg(k, 0, sizeof( gOutBuffer[gMinVectorSizeIndex]), &gOutBuffer[gMinVectorSizeIndex]))) + if ((error = + clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]), + &gOutBuffer[gMinVectorSizeIndex]))) { vlog_error( "Error: Unable to set kernel arg to detect RTZ mode for the device. Err = %d", error ); return error; } size_t dim = 1; - if((error = clEnqueueNDRangeKernel(gQueue, k, 1, NULL, &dim, NULL, 0, NULL, NULL) )) + if ((error = clEnqueueNDRangeKernel(gQueue, kernel, 1, NULL, &dim, NULL, 0, + NULL, NULL))) { vlog_error( "Error: Unable to execute kernel to detect RTZ mode for the device. Err = %d", error ); return error; @@ -1623,9 +1608,6 @@ static int IsInRTZMode( void ) return error; } - clReleaseKernel(k); - clReleaseProgram(query); - return data.isRTZ; } diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index b169e6b9..2b804e40 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -306,15 +306,22 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont if(allTestCase[testId]->_type == TYPE_VECTOR) { - err = create_single_kernel_helper(context, &program, NULL, sizeof(sourceVec) / sizeof(sourceVec[0]), sourceVec, NULL); + err = create_single_kernel_helper( + context, &program, kernel_ptr, + sizeof(sourceVec) / sizeof(sourceVec[0]), sourceVec, testname); } else if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { - err = create_single_kernel_helper(context, &program, NULL, sizeof(sourceAddrSpace) / sizeof(sourceAddrSpace[0]), sourceAddrSpace, NULL); + err = create_single_kernel_helper(context, &program, kernel_ptr, + sizeof(sourceAddrSpace) + / sizeof(sourceAddrSpace[0]), + sourceAddrSpace, testname); } else { - err = create_single_kernel_helper(context, &program, NULL, sizeof(sourceGen) / sizeof(sourceGen[0]), sourceGen, NULL); + err = create_single_kernel_helper( + context, &program, kernel_ptr, + sizeof(sourceGen) / sizeof(sourceGen[0]), sourceGen, testname); } if (!program || err) { @@ -322,12 +329,6 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont return NULL; } - *kernel_ptr = clCreateKernel(program, testname, &err); - if ( err ) { - log_error("clCreateKernel failed (%d)\n", err); - return NULL; - } - return program; } diff --git a/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp b/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp index 84f8ed1f..9e1789c2 100644 --- a/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp +++ b/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp @@ -129,13 +129,9 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID, { // Run the cl kernel for reference results clProgramWrapper prog; - err = create_single_kernel_helper_create_program(context, &prog, 1, &kernelBuf, NULL); - SPIRV_CHECK_ERROR(err, "Failed to create cl program"); - - err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to build program"); - - clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err); + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &prog, &kernel, 1, + &kernelBuf, "fmath_cl"); SPIRV_CHECK_ERROR(err, "Failed to create cl kernel"); clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); diff --git a/test_conformance/spirv_new/test_op_fmath.cpp b/test_conformance/spirv_new/test_op_fmath.cpp index 7250eb15..bec0667c 100644 --- a/test_conformance/spirv_new/test_op_fmath.cpp +++ b/test_conformance/spirv_new/test_op_fmath.cpp @@ -89,13 +89,9 @@ int test_fmath(cl_device_id deviceID, { // Run the cl kernel for reference results clProgramWrapper prog; - err = create_single_kernel_helper_create_program(context, &prog, 1, &kernelBuf, NULL); - SPIRV_CHECK_ERROR(err, "Failed to create cl program"); - - err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to build program"); - - clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err); + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &prog, &kernel, 1, + &kernelBuf, "fmath_cl"); SPIRV_CHECK_ERROR(err, "Failed to create cl kernel"); clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); diff --git a/test_conformance/spirv_new/test_op_vector_times_scalar.cpp b/test_conformance/spirv_new/test_op_vector_times_scalar.cpp index 99d71f72..0a604bcf 100644 --- a/test_conformance/spirv_new/test_op_vector_times_scalar.cpp +++ b/test_conformance/spirv_new/test_op_vector_times_scalar.cpp @@ -82,15 +82,11 @@ int test_vector_times_scalar(cl_device_id deviceID, { // Run the cl kernel for reference results clProgramWrapper prog; - err = create_single_kernel_helper_create_program(context, &prog, 1, &kernelBuf, NULL); + clKernelWrapper kernel; + err = create_single_kernel_helper(context, &prog, &kernel, 1, + &kernelBuf, "vector_times_scalar"); SPIRV_CHECK_ERROR(err, "Failed to create cl program"); - err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL); - SPIRV_CHECK_ERROR(err, "Failed to build program"); - - clKernelWrapper kernel = clCreateKernel(prog, "vector_times_scalar", &err); - SPIRV_CHECK_ERROR(err, "Failed to create cl kernel"); - clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, res_bytes, NULL, &err); SPIRV_CHECK_ERROR(err, "Failed to create ref buffer");