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 <ellen.norris-thompson@arm.com>

* Skip tests using clCompileProgram in offline mode

Contributes #31

Signed-off-by: Ellen Norris-Thompson <ellen.norris-thompson@arm.com>

* Using type wrappers when using kernel helper functions

Also includes fix for windows build

Fixes #31

Signed-off-by: Ellen Norris-Thompson <ellen.norris-thompson@arm.com>

* Remove clReleaseKernel for wrapped kernel

Fixes #31

Signed-off-by: Ellen Norris-Thompson <ellen.norris-thompson@arm.com>

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,

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,
+        clKernelWrapper kernel;
-                                                         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);
+        err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &inBuf);
-        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);
         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,
+        err = clEnqueueNDRangeKernel(q, kernel, 1, NULL, &global_work_size,
-                                     NULL, NULL);
+                                     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;
     }
 

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);
+        error = create_single_kernel_helper(
-        test_error( error, "Unable to build test program" );
+            local_context, &local_program, &local_kernel, 1,
-
+            &repeate_test_kernel, "test_kernel");
-        local_kernel = clCreateKernel(local_program, "test_kernel", &error);
+        test_error(error, "Unable to create kernel");
-        test_error( error, "clCreateKernel failed");
 
         local_mem_in = clCreateBuffer(local_context, CL_MEM_READ_ONLY, TEST_SIZE*sizeof(cl_int), NULL, &error);
         test_error( error, "clCreateBuffer failed");

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");
+    test_error(status, "Unable to create kernel");
-
-    kernel = clCreateKernel(program, "test_kernel", &status);
-    test_error(status, "CreateKernel failed.");
 
     cl_mem dev_src = clCreateBuffer(context, CL_MEM_READ_ONLY, NITEMS*sizeof(cl_float),
         NULL, NULL);

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 );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
-    test_error( err, "Unable to build sample program" );
+                                      testProgram, "sample_test");
-
+    test_error(err, "Unable to build sample program and sample_test kernel");
-    kernel = clCreateKernel( program, "sample_test", &err );
-    test_error( err, "Unable to create sample_test kernel" );
 
     err = clSetKernelArg( kernel, 0, sizeof(cl_mem), &buffer );
     test_error( err, "Unable to set kernel argument" );

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 */
+    /* Create a test program and kernel from it */
-    error = create_single_kernel_helper(context, &program, NULL, 1, testProgram, NULL);
+    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 );

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;

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]);
+        status = create_single_kernel_helper(context, &program, &kernel, 1,
-        TestRequire(
+                                             &sourceTexts[0], entrypoint);
-            CL_SUCCESS == status,
+        TestRequire(CL_SUCCESS == status, "Kernel creation failed");
-            "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;

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 <sys/sysctl.h>
@@ -1384,36 +1385,36 @@ void _LogBuildError( cl_program p, int line, const char *file )
 int InitILogbConstants( void )
 {
     int error;
-    const char *kernel =
+    const char *kernelSource =
-    "__kernel void GetILogBConstants( __global int *out )\n"
+        R"(__kernel void GetILogBConstants( __global int *out )
-    "{\n"
+        {
-    "   out[0] = FP_ILOGB0;\n"
+            out[0] = FP_ILOGB0;
-    "   out[1] = FP_ILOGBNAN;\n"
+            out[1] = FP_ILOGBNAN;
-    "}\n";
+        })";
 
-    cl_program query;
+    clProgramWrapper query;
-    error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL);
+    clKernelWrapper kernel;
-    if (NULL == query || error)
+    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 ((error =
-    if( NULL == k || error)
+             clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]),
-    {
+                            &gOutBuffer[gMinVectorSizeIndex])))
-      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])))
     {
         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 =
+    const char *kernelSource =
-    "__kernel void IsTininessDetectedBeforeRounding( __global float *out )\n"
+        R"(__kernel void IsTininessDetectedBeforeRounding( __global float *out )
-    "{\n"
+        {
-    "   volatile float a = 0x1.000002p-126f;\n"
+           volatile float a = 0x1.000002p-126f;
-    "   volatile float b = 0x1.fffffcp-1f;\n"       // product is 0x1.fffffffffff8p-127
+           volatile float b = 0x1.fffffcp-1f;
-    "   out[0] = a * b;\n"
+           out[0] = a * b; // product is 0x1.fffffffffff8p-127
-    "}\n";
+        })";
 
-    cl_program query;
+    clProgramWrapper query;
-    error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL);
+    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 ((error =
-    if( NULL == k || error)
+             clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]),
-    {
+                            &gOutBuffer[gMinVectorSizeIndex])))
-      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])))
     {
         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);
+        vlog_error("\t\tFAILED -- Failed to create kernel. (%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 );
         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 =
+    const char *kernelSource =
-    "__kernel void GetRoundingMode( __global int *out )\n"
+        R"(__kernel void GetRoundingMode( __global int *out )
-    "{\n"
+        {
-    "   volatile float a = 0x1.0p23f;\n"
+            volatile float a = 0x1.0p23f;
-    "   volatile float b = -0x1.0p23f;\n"
+            volatile float b = -0x1.0p23f;
-    "   out[0] = (a + 0x1.fffffep-1f == a) && (b - 0x1.fffffep-1f == b);\n"
+            out[0] = (a + 0x1.fffffep-1f == a) && (b - 0x1.fffffep-1f == b);
-    "}\n";
+        "})";
 
-    cl_program query;
+    clProgramWrapper query;
-    error = create_single_kernel_helper(gContext, &query, NULL, 1, &kernel, NULL);
+    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 ((error =
-    if( NULL == k || error)
+             clSetKernelArg(kernel, 0, sizeof(gOutBuffer[gMinVectorSizeIndex]),
-    {
+                            &gOutBuffer[gMinVectorSizeIndex])))
-        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])))
     {
         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;
 }
 

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;
 }
 

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);
+        clKernelWrapper kernel;
-        SPIRV_CHECK_ERROR(err, "Failed to create cl program");
+        err = create_single_kernel_helper(context, &prog, &kernel, 1,
-
+                                          &kernelBuf, "fmath_cl");
-        err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL);
-        SPIRV_CHECK_ERROR(err, "Failed to build program");
-
-        clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err);
         SPIRV_CHECK_ERROR(err, "Failed to create cl kernel");
 
         clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);

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);
+        clKernelWrapper kernel;
-        SPIRV_CHECK_ERROR(err, "Failed to create cl program");
+        err = create_single_kernel_helper(context, &prog, &kernel, 1,
-
+                                          &kernelBuf, "fmath_cl");
-        err = clBuildProgram(prog, 1, &deviceID, NULL, NULL, NULL);
-        SPIRV_CHECK_ERROR(err, "Failed to build program");
-
-        clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err);
         SPIRV_CHECK_ERROR(err, "Failed to create cl kernel");
 
         clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);

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");