Use highest OpenCL C version (#1081)

* Set the highest supported OpenCL C version.

* Remove gDeviceLt20 variable - not used anymore.

* Fix formatting issues

test_conformance/SVM/main.cpp

@@ -217,7 +217,8 @@ cl_int create_cl_objects(cl_device_id device_from_harness, const char** ppCodeSt
     {
         if (!is_extension_available(devices[i], extension.c_str()))
         {
-        log_error("Required extension not found - device id %d - %s\n", i, extension.c_str());
+            log_error("Required extension not found - device id %d - %s\n", i,
+                      extension.c_str());
             extensions_supported = false;
             break;
         }
@@ -249,10 +250,11 @@ cl_int create_cl_objects(cl_device_id device_from_harness, const char** ppCodeSt
     test_error(error, "clCreateCommandQueue failed");
   }
 
-  if(ppCodeString)
+  if (ppCodeString)
   {
-    error = create_single_kernel_helper(*context, program, 0, 1, ppCodeString, 0, "-cl-std=CL2.0");
-    test_error( error, "failed to create program" );
+      error =
+          create_single_kernel_helper(*context, program, 0, 1, ppCodeString, 0);
+      test_error(error, "failed to create program");
   }
 
   return 0;

test_conformance/api/test_sub_group_dispatch.cpp

@@ -95,7 +95,9 @@ int test_sub_group_dispatch(cl_device_id deviceID, cl_context context, cl_comman
         }
     }
 
-    error = create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, subgroup_dispatch_kernel, "subgroup_dispatch_kernel", "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        subgroup_dispatch_kernel,
+                                        "subgroup_dispatch_kernel");
     if (error != 0)
         return error;
 

test_conformance/device_execution/enqueue_profiling.cpp

@@ -89,9 +89,9 @@ int test_enqueue_profiling(cl_device_id device, cl_context context,
 
     cl_event kernel_event;
 
-    err_ret = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel, 1, &enqueue_multi_level,
-        "enqueue_multi_level", "-cl-std=CL2.0");
+    err_ret = create_single_kernel_helper(context, &program, &kernel, 1,
+                                          &enqueue_multi_level,
+                                          "enqueue_multi_level");
     if (check_error(err_ret, "Create single kernel failed")) return -1;
 
     res_mem = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,

test_conformance/device_execution/host_multi_queue.cpp

@@ -184,7 +184,11 @@ int test_host_multi_queue(cl_device_id device, cl_context context, cl_command_qu
                 global = 16;
             }
 
-            err_ret |= create_single_kernel_helper_with_build_options(context, &program[i], &kernel[i], sources_multi_queue_block[i].num_lines, sources_multi_queue_block[i].lines, sources_multi_queue_block[i].kernel_name, "-cl-std=CL2.0");
+            err_ret |= create_single_kernel_helper(
+                context, &program[i], &kernel[i],
+                sources_multi_queue_block[i].num_lines,
+                sources_multi_queue_block[i].lines,
+                sources_multi_queue_block[i].kernel_name);
             if(check_error(err_ret, "Create single kernel failed")) { res = -1; break; }
 
             mem[i] = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(kernel_results), kernel_results, &err_ret);

test_conformance/device_execution/host_queue_order.cpp

@@ -133,10 +133,14 @@ int test_host_queue_order(cl_device_id device, cl_context context, cl_command_qu
 
     cl_event kernel_event;
 
-    err_ret = create_single_kernel_helper_with_build_options(context, &program1, &kernel1,  arr_size(enqueue_block_first_kernel), enqueue_block_first_kernel, "enqueue_block_first_kernel", "-cl-std=CL2.0");
+    err_ret = create_single_kernel_helper(
+        context, &program1, &kernel1, arr_size(enqueue_block_first_kernel),
+        enqueue_block_first_kernel, "enqueue_block_first_kernel");
     if(check_error(err_ret, "Create single kernel failed")) return -1;
 
-    err_ret = create_single_kernel_helper_with_build_options(context, &program2, &kernel2, arr_size(enqueue_block_second_kernel), enqueue_block_second_kernel, "enqueue_block_second_kernel", "-cl-std=CL2.0");
+    err_ret = create_single_kernel_helper(
+        context, &program2, &kernel2, arr_size(enqueue_block_second_kernel),
+        enqueue_block_second_kernel, "enqueue_block_second_kernel");
     if(check_error(err_ret, "Create single kernel failed")) return -1;
 
     res_mem = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, sizeof(kernel_results), kernel_results, &err_ret);

test_conformance/device_execution/utils.cpp

@@ -40,7 +40,8 @@ int run_n_kernel_args(cl_context context, cl_command_queue queue, const char** s
     cl_uint i;
     size_t ret_len;
 
-    err_ret = create_single_kernel_helper_with_build_options(context, &program, &kernel, num_lines, source, kernel_name, "-cl-std=CL2.0");
+    err_ret = create_single_kernel_helper(context, &program, &kernel, num_lines,
+                                          source, kernel_name);
     if(check_error(err_ret, "Create single kernel failed")) return -1;
 
     mem = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, res_size, results, &err_ret);

test_conformance/generic_address_space/basic_tests.cpp

@@ -39,7 +39,9 @@ public:
 
         const char *srcPtr = src.c_str();
 
-        if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) {
+        if (create_single_kernel_helper(context, &program, &kernel, 1, &srcPtr,
+                                        "testKernel"))
+        {
             log_error("create_single_kernel_helper failed");
             return -1;
         }

test_conformance/generic_address_space/stress_tests.cpp

@@ -41,7 +41,9 @@ public:
 
         const char *srcPtr = src.c_str();
 
-        if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) {
+        if (create_single_kernel_helper(context, &program, &kernel, 1, &srcPtr,
+                                        "testKernel"))
+        {
             log_error("create_single_kernel_helper failed");
             return -1;
         }

test_conformance/images/kernel_image_methods/main.cpp

@@ -23,7 +23,6 @@
 bool gDebugTrace;
 bool gTestSmallImages;
 bool gTestMaxImages;
-bool gDeviceLt20 = false;
 
 cl_channel_type gChannelTypeToUse = (cl_channel_type)-1;
 cl_channel_order gChannelOrderToUse = (cl_channel_order)-1;

test_conformance/images/kernel_image_methods/test_1D.cpp

@@ -15,7 +15,6 @@
 //
 #include "../testBase.h"
 
-extern bool gDeviceLt20;
 
 struct image_kernel_data
 {
@@ -98,7 +97,8 @@ static int test_get_1Dimage_info_single(cl_context context,
     if (error)
         print_error(error, "clFinish failed.\n");
     const char *ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create kernel to test against" );
 
     // Create an output buffer

test_conformance/images/kernel_image_methods/test_1D_array.cpp

@@ -15,7 +15,6 @@
 //
 #include "../testBase.h"
 
-extern bool gDeviceLt20;
 
 struct image_kernel_data
 {
@@ -102,7 +101,8 @@ int test_get_1Dimage_array_info_single(cl_context context,
     if (error)
         print_error(error, "clFinish failed.\n");
     const char *ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create kernel to test against" );
 
     // Create an output buffer

test_conformance/images/kernel_image_methods/test_2D.cpp

@@ -15,7 +15,6 @@
 //
 #include "../testBase.h"
 
-extern bool gDeviceLt20;
 
 struct image_kernel_data
 {
@@ -133,7 +132,8 @@ int test_get_image_info_single(cl_context context, cl_command_queue queue,
     if (error)
         print_error(error, "clFinish failed.\n");
     const char *ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create kernel to test against" );
 
     // Create an output buffer

test_conformance/images/kernel_image_methods/test_2D_array.cpp

@@ -15,7 +15,6 @@
 //
 #include "../testBase.h"
 
-extern bool gDeviceLt20;
 
 struct image_kernel_data
 {
@@ -108,7 +107,8 @@ int test_get_2Dimage_array_info_single(cl_context context,
     if (error)
         print_error(error, "clFinish failed.\n");
     const char *ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create kernel to test against" );
 
     // Create an output buffer

test_conformance/images/kernel_image_methods/test_loops.cpp

@@ -16,7 +16,6 @@
 #include "../testBase.h"
 #include "../common.h"
 
-extern bool gDeviceLt20;
 
 extern int test_get_image_info_1D(cl_device_id device, cl_context context,
                                   cl_command_queue queue,
@@ -117,9 +116,6 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu
 {
     int version_check;
     auto version = get_device_cl_version(device);
-    if (version < Version(2, 0)) {
-        gDeviceLt20 = true;
-    }
 
     if ((version_check = (version < Version(1, 2))))
     {

test_conformance/images/kernel_read_write/main.cpp

@@ -35,7 +35,6 @@ bool gTestSmallImages;
 bool gTestMaxImages;
 bool gTestImage2DFromBuffer;
 bool gTestMipmaps;
-bool gDeviceLt20 = false;
 cl_filter_mode    gFilterModeToUse = (cl_filter_mode)-1;
 // Default is CL_MEM_USE_HOST_PTR for the test
 cl_mem_flags    gMemFlagsToUse = CL_MEM_USE_HOST_PTR;
@@ -107,10 +106,6 @@ static int doTest( cl_device_id device, cl_context context, cl_command_queue que
     bool            tDisableOffsets = false;
     bool            tNormalizedModeToUse = false;
     cl_filter_mode  tFilterModeToUse = (cl_filter_mode)-1;
-    auto version = get_device_cl_version(device);
-    if (version < Version(2, 0)) {
-        gDeviceLt20 = true;
-    }
 
     if( testTypesToRun & kReadTests )
     {

test_conformance/images/kernel_read_write/test_iterations.cpp

@@ -26,7 +26,6 @@ extern bool gTestImage2DFromBuffer;
 extern uint64_t gRoundingStartValue;
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 // Utility function to clamp down image sizes for certain tests to avoid
 // using too much memory.
@@ -1664,7 +1663,8 @@ int test_read_image_set_2D( cl_device_id device, cl_context context, cl_command_
             gTestMipmaps?", lod":" ");
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if( gTestSmallImages )

test_conformance/images/kernel_read_write/test_read_1D.cpp

@@ -26,7 +26,6 @@
 extern uint64_t gRoundingStartValue;
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 const char *read1DKernelSourcePattern =
 "__kernel void sample_kernel( read_only image1d_t input,%s __global float *xOffsets, __global %s4 *results %s)\n"
@@ -1056,7 +1055,8 @@ int test_read_image_set_1D( cl_device_id device, cl_context context, cl_command_
 
     ptr = programSrc;
 
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     if(error)
     {
         exit(1);

test_conformance/images/kernel_read_write/test_read_1D_array.cpp

@@ -25,7 +25,6 @@
 extern uint64_t gRoundingStartValue;
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 const char *read1DArrayKernelSourcePattern =
 "__kernel void sample_kernel( read_only image1d_array_t input,%s __global float *xOffsets, __global float *yOffsets, __global %s4 *results %s)\n"
@@ -1165,7 +1164,8 @@ int test_read_image_set_1D_array( cl_device_id device, cl_context context, cl_co
             gTestMipmaps ? ", lod" : "" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if( gTestSmallImages )

test_conformance/images/kernel_read_write/test_read_2D_array.cpp

@@ -18,7 +18,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 // Utility function to clamp down image sizes for certain tests to avoid
 // using too much memory.
@@ -1392,7 +1391,8 @@ int test_read_image_set_2D_array( cl_device_id device, cl_context context, cl_co
             gTestMipmaps ? ", lod" : " " );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_read_3D.cpp

@@ -18,7 +18,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 // Utility function to clamp down image sizes for certain tests to avoid
 // using too much memory.
@@ -1230,7 +1229,8 @@ int test_read_image_set_3D( cl_device_id device, cl_context context, cl_command_
             gTestMipmaps? ",lod":" ");
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_write_1D.cpp

@@ -21,7 +21,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor* imageInfo );
@@ -580,7 +579,8 @@ int test_write_image_1D_set( cl_device_id device, cl_context context, cl_command
              gTestMipmaps ? ", lod" :"" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_write_1D_array.cpp

@@ -21,7 +21,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
@@ -603,7 +602,8 @@ int test_write_image_1D_array_set( cl_device_id device, cl_context context, cl_c
              gTestMipmaps ? ", lod" :"" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_write_2D_array.cpp

@@ -21,7 +21,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
@@ -630,7 +629,8 @@ int test_write_image_2D_array_set( cl_device_id device, cl_context context, cl_c
              gTestMipmaps ? ", lod" : "" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_write_3D.cpp

@@ -21,7 +21,6 @@
 
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
@@ -636,7 +635,8 @@ int test_write_image_3D_set( cl_device_id device, cl_context context, cl_command
              gTestMipmaps ? ", lod" : "" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/kernel_read_write/test_write_image.cpp

@@ -22,7 +22,6 @@
 extern bool gTestImage2DFromBuffer;
 extern cl_mem_flags gMemFlagsToUse;
 extern int gtestTypesToRun;
-extern bool gDeviceLt20;
 
 extern int test_write_image_1D_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d );
 extern int test_write_image_3D_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d );
@@ -682,7 +681,8 @@ int test_write_image_set( cl_device_id device, cl_context context, cl_command_qu
              gTestMipmaps ? ", lod" : "" );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     // Run tests

test_conformance/images/samplerlessReads/main.cpp

@@ -36,7 +36,6 @@ int                 gTypesToTest;
 cl_channel_type     gChannelTypeToUse = (cl_channel_type)-1;
 cl_channel_order    gChannelOrderToUse = (cl_channel_order)-1;
 bool                gEnablePitch = false;
-bool                gDeviceLt20 = false;
 
 static void printUsage( const char *execName );
 

test_conformance/images/samplerlessReads/test_iterations.cpp

@@ -22,7 +22,6 @@
     #include <setjmp.h>
 #endif
 
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 const char *read2DKernelSourcePattern =
@@ -254,7 +253,8 @@ int test_read_image_set_2D( cl_device_id device, cl_context context, cl_command_
     }
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if ( gTestSmallImages )

test_conformance/images/samplerlessReads/test_loops.cpp

@@ -17,7 +17,6 @@
 #include "../common.h"
 
 extern int gTypesToTest;
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 extern int test_read_image_set_1D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, image_sampler_data *imageSampler, ExplicitType outputType );
@@ -99,10 +98,6 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu
     cl_image_format *formatList;
     bool *filterFlags;
     unsigned int numFormats;
-    auto version = get_device_cl_version(device);
-    if (version < Version(2, 0)) {
-        gDeviceLt20 = true;
-    }
 
     if (gTestReadWrite && checkForReadWriteImageSupport(device))
     {

test_conformance/images/samplerlessReads/test_read_1D.cpp

@@ -22,7 +22,6 @@
     #include <setjmp.h>
 #endif
 
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 const char *read1DKernelSourcePattern =
@@ -252,7 +251,8 @@ int test_read_image_set_1D( cl_device_id device, cl_context context, cl_command_
 
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if ( gTestSmallImages )

test_conformance/images/samplerlessReads/test_read_1D_array.cpp

@@ -22,7 +22,6 @@
     #include <setjmp.h>
 #endif
 
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 const char *read1DArrayKernelSourcePattern =
@@ -251,7 +250,8 @@ int test_read_image_set_1D_array( cl_device_id device, cl_context context, cl_co
 
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if ( gTestSmallImages )

test_conformance/images/samplerlessReads/test_read_1D_buffer.cpp

@@ -22,7 +22,6 @@
     #include <setjmp.h>
 #endif
 
-extern bool gDeviceLt20;
 
 const char *read1DBufferKernelSourcePattern =
 "__kernel void sample_kernel( read_only image1d_buffer_t inputA, read_only image1d_t inputB, sampler_t sampler, __global int *results )\n"
@@ -244,7 +243,8 @@ int test_read_image_set_1D_buffer( cl_device_id device, cl_context context, cl_c
              readFormat );
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
     if ( gTestSmallImages )

test_conformance/images/samplerlessReads/test_read_2D_array.cpp

@@ -16,7 +16,6 @@
 #include "../testBase.h"
 #include <float.h>
 
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 const char *read2DArrayKernelSourcePattern =
@@ -241,7 +240,8 @@ int test_read_image_set_2D_array( cl_device_id device, cl_context context, cl_co
 
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
 

test_conformance/images/samplerlessReads/test_read_3D.cpp

@@ -16,7 +16,6 @@
 #include "../testBase.h"
 #include <float.h>
 
-extern bool gDeviceLt20;
 extern bool gTestReadWrite;
 
 const char *read3DKernelSourcePattern =
@@ -244,7 +243,8 @@ int test_read_image_set_3D( cl_device_id device, cl_context context, cl_command_
 
 
     ptr = programSrc;
-    error = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &ptr, "sample_kernel", gDeviceLt20 ? "" : "-cl-std=CL2.0" );
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
     test_error( error, "Unable to create testing kernel" );
 
 

test_conformance/pipes/test_pipe_limits.cpp

@@ -218,9 +218,8 @@ int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_que
     const char *sources[] = { kernel_source.c_str() };
 
     // Create producer kernel
-    err = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel[0], 1, sources, kernelName[0],
-        "-cl-std=CL2.0");
+    err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
+                                      kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel
@@ -368,9 +367,8 @@ int test_pipe_max_packet_size(cl_device_id deviceID, cl_context context, cl_comm
     const char *sources[] = { kernel_source.c_str() };
 
     // Create producer kernel
-    err = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel[0], 1, sources, kernelName[0],
-        "-cl-std=CL2.0");
+    err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
+                                      kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel
@@ -533,9 +531,8 @@ int test_pipe_max_active_reservations(cl_device_id deviceID, cl_context context,
     const char *sources[] = { kernel_source.c_str() };
 
     // Create producer kernel
-    err = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel[0], 1, sources, kernelName[0],
-        "-cl-std=CL2.0");
+    err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
+                                      kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     // Create consumer kernel

test_conformance/pipes/test_pipe_query_functions.cpp

@@ -140,7 +140,9 @@ int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_comm
     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");
+    err = create_single_kernel_helper(
+        context, &program, &kernel[0], 1,
+        (const char **)&pipe_query_functions_kernel_code, kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel

test_conformance/pipes/test_pipe_read_write.cpp

@@ -531,9 +531,8 @@ int test_pipe_readwrite( cl_device_id deviceID, cl_context context, cl_command_q
         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, sources, kernelName[ii],
-            "-cl-std=CL2.0");
+        err = create_single_kernel_helper(context, &program[i], &kernel[ii], 1,
+                                          sources, kernelName[ii]);
 
         test_error_ret(err, " Error creating program", -1);
 
@@ -659,7 +658,8 @@ int test_pipe_readwrite_struct_generic( cl_device_id deviceID, cl_context contex
     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");
+    err = create_single_kernel_helper(context, &program, &kernel[0], 1,
+                                      &kernelCode, kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel

test_conformance/pipes/test_pipe_readwrite_errors.cpp

@@ -115,7 +115,9 @@ int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_com
     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");
+    err = create_single_kernel_helper(
+        context, &program, &kernel[0], 1,
+        (const char **)&pipe_readwrite_errors_kernel_code, kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel

test_conformance/pipes/test_pipe_subgroups.cpp

@@ -146,7 +146,9 @@ int test_pipe_subgroups_divergence(cl_device_id deviceID, cl_context context, cl
     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");
+    err = create_single_kernel_helper(
+        context, &program, &kernel[0], 1,
+        (const char **)&pipe_subgroups_kernel_code, kernelName[0]);
     test_error_ret(err, " Error creating program", -1);
 
     //Create consumer kernel

test_conformance/subgroups/subhelpers.h

@@ -377,8 +377,8 @@ struct test
         const std::string &kernel_str = kernel_sstr.str();
         const char *kernel_src = kernel_str.c_str();
 
-        error = create_single_kernel_helper_with_build_options(
-            context, &program, &kernel, 1, &kernel_src, kname, "-cl-std=CL2.0");
+        error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                            &kernel_src, kname);
         if (error != 0) return error;
 
         // Determine some local dimensions to use for the test.

test_conformance/subgroups/test_queries.cpp

@@ -67,9 +67,8 @@ int test_sub_group_info(cl_device_id device, cl_context context,
 
     const std::string &kernel_str = kernel_sstr.str();
     const char *kernel_src = kernel_str.c_str();
-    error = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel, 1, &kernel_src, "query_kernel",
-        "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &kernel_src, "query_kernel");
     if (error != 0) return error;
 
     // Determine some local dimensions to use for the test.

test_conformance/subgroups/test_workitem.cpp

@@ -227,9 +227,8 @@ int test_work_item_functions(cl_device_id device, cl_context context,
            "}";
     const std::string &kernel_str = kernel_sstr.str();
     const char *kernel_src = kernel_str.c_str();
-    error = create_single_kernel_helper_with_build_options(
-        context, &program, &kernel, 1, &kernel_src, "get_test",
-        "-cl-std=CL2.0");
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &kernel_src, "get_test");
     if (error != 0) return error;
 
     error = get_max_allowed_work_group_size(context, kernel, &local, NULL);

test_conformance/workgroups/test_wg_all.cpp

@@ -79,7 +79,8 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_all_kernel_code, "test_wg_all", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_all_kernel_code, "test_wg_all");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_any.cpp

@@ -79,7 +79,8 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_any_kernel_code, "test_wg_any", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_any_kernel_code, "test_wg_any");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_broadcast.cpp

@@ -174,7 +174,9 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_broadcast_1D_kernel_code, "test_wg_broadcast_1D", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_broadcast_1D_kernel_code,
+                                      "test_wg_broadcast_1D");
     if (err)
         return -1;
 
@@ -281,7 +283,9 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_broadcast_2D_kernel_code, "test_wg_broadcast_2D", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_broadcast_2D_kernel_code,
+                                      "test_wg_broadcast_2D");
     if (err)
         return -1;
 
@@ -406,7 +410,9 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_broadcast_3D_kernel_code, "test_wg_broadcast_3D", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_broadcast_3D_kernel_code,
+                                      "test_wg_broadcast_3D");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_reduce.cpp

@@ -176,7 +176,9 @@ test_work_group_reduce_add_int(cl_device_id device, cl_context context, cl_comma
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_add_kernel_code_int, "test_wg_reduce_add_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_add_kernel_code_int,
+                                      "test_wg_reduce_add_int");
     if (err)
         return -1;
 
@@ -279,7 +281,9 @@ test_work_group_reduce_add_uint(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_add_kernel_code_uint, "test_wg_reduce_add_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_add_kernel_code_uint,
+                                      "test_wg_reduce_add_uint");
     if (err)
         return -1;
 
@@ -381,7 +385,9 @@ test_work_group_reduce_add_long(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_add_kernel_code_long, "test_wg_reduce_add_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_add_kernel_code_long,
+                                      "test_wg_reduce_add_long");
     if (err)
         return -1;
 
@@ -484,7 +490,9 @@ test_work_group_reduce_add_ulong(cl_device_id device, cl_context context, cl_com
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_add_kernel_code_ulong, "test_wg_reduce_add_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_add_kernel_code_ulong,
+                                      "test_wg_reduce_add_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_reduce_max.cpp

@@ -177,7 +177,9 @@ test_work_group_reduce_max_int(cl_device_id device, cl_context context, cl_comma
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_max_kernel_code_int, "test_wg_reduce_max_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_max_kernel_code_int,
+                                      "test_wg_reduce_max_int");
     if (err)
         return -1;
 
@@ -289,7 +291,9 @@ test_work_group_reduce_max_uint(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_max_kernel_code_uint, "test_wg_reduce_max_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_max_kernel_code_uint,
+                                      "test_wg_reduce_max_uint");
     if (err)
         return -1;
 
@@ -400,7 +404,9 @@ test_work_group_reduce_max_long(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_max_kernel_code_long, "test_wg_reduce_max_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_max_kernel_code_long,
+                                      "test_wg_reduce_max_long");
     if (err)
         return -1;
 
@@ -512,7 +518,9 @@ test_work_group_reduce_max_ulong(cl_device_id device, cl_context context, cl_com
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_max_kernel_code_ulong, "test_wg_reduce_max_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_max_kernel_code_ulong,
+                                      "test_wg_reduce_max_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_reduce_min.cpp

@@ -177,7 +177,9 @@ test_work_group_reduce_min_int(cl_device_id device, cl_context context, cl_comma
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_min_kernel_code_int, "test_wg_reduce_min_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_min_kernel_code_int,
+                                      "test_wg_reduce_min_int");
     if (err)
         return -1;
 
@@ -289,7 +291,9 @@ test_work_group_reduce_min_uint(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_min_kernel_code_uint, "test_wg_reduce_min_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_min_kernel_code_uint,
+                                      "test_wg_reduce_min_uint");
     if (err)
         return -1;
 
@@ -400,7 +404,9 @@ test_work_group_reduce_min_long(cl_device_id device, cl_context context, cl_comm
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_min_kernel_code_long, "test_wg_reduce_min_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_min_kernel_code_long,
+                                      "test_wg_reduce_min_long");
     if (err)
         return -1;
 
@@ -512,7 +518,9 @@ test_work_group_reduce_min_ulong(cl_device_id device, cl_context context, cl_com
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_reduce_min_kernel_code_ulong, "test_wg_reduce_min_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_reduce_min_kernel_code_ulong,
+                                      "test_wg_reduce_min_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_exclusive_add.cpp

@@ -184,7 +184,9 @@ test_work_group_scan_exclusive_add_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_add_kernel_code_int, "test_wg_scan_exclusive_add_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_add_kernel_code_int,
+                                      "test_wg_scan_exclusive_add_int");
     if (err)
         return -1;
 
@@ -287,7 +289,9 @@ test_work_group_scan_exclusive_add_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_add_kernel_code_uint, "test_wg_scan_exclusive_add_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_add_kernel_code_uint,
+                                      "test_wg_scan_exclusive_add_uint");
     if (err)
         return -1;
 
@@ -389,7 +393,9 @@ test_work_group_scan_exclusive_add_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_add_kernel_code_long, "test_wg_scan_exclusive_add_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_add_kernel_code_long,
+                                      "test_wg_scan_exclusive_add_long");
     if (err)
         return -1;
 
@@ -492,7 +498,9 @@ test_work_group_scan_exclusive_add_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_add_kernel_code_ulong, "test_wg_scan_exclusive_add_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_add_kernel_code_ulong,
+                                      "test_wg_scan_exclusive_add_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_exclusive_max.cpp

@@ -176,7 +176,9 @@ test_work_group_scan_exclusive_max_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_max_kernel_code_int, "test_wg_scan_exclusive_max_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_max_kernel_code_int,
+                                      "test_wg_scan_exclusive_max_int");
     if (err)
         return -1;
 
@@ -288,7 +290,9 @@ test_work_group_scan_exclusive_max_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_max_kernel_code_uint, "test_wg_scan_exclusive_max_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_max_kernel_code_uint,
+                                      "test_wg_scan_exclusive_max_uint");
     if (err)
         return -1;
 
@@ -399,7 +403,9 @@ test_work_group_scan_exclusive_max_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_max_kernel_code_long, "test_wg_scan_exclusive_max_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_max_kernel_code_long,
+                                      "test_wg_scan_exclusive_max_long");
     if (err)
         return -1;
 
@@ -511,7 +517,9 @@ test_work_group_scan_exclusive_max_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_max_kernel_code_ulong, "test_wg_scan_exclusive_max_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_max_kernel_code_ulong,
+                                      "test_wg_scan_exclusive_max_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_exclusive_min.cpp

@@ -177,7 +177,9 @@ test_work_group_scan_exclusive_min_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_min_kernel_code_int, "test_wg_scan_exclusive_min_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_min_kernel_code_int,
+                                      "test_wg_scan_exclusive_min_int");
     if (err)
         return -1;
 
@@ -289,7 +291,9 @@ test_work_group_scan_exclusive_min_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_min_kernel_code_uint, "test_wg_scan_exclusive_min_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_min_kernel_code_uint,
+                                      "test_wg_scan_exclusive_min_uint");
     if (err)
         return -1;
 
@@ -400,7 +404,9 @@ test_work_group_scan_exclusive_min_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_min_kernel_code_long, "test_wg_scan_exclusive_min_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_min_kernel_code_long,
+                                      "test_wg_scan_exclusive_min_long");
     if (err)
         return -1;
 
@@ -512,7 +518,9 @@ test_work_group_scan_exclusive_min_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_exclusive_min_kernel_code_ulong, "test_wg_scan_exclusive_min_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_exclusive_min_kernel_code_ulong,
+                                      "test_wg_scan_exclusive_min_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_inclusive_add.cpp

@@ -173,7 +173,9 @@ test_work_group_scan_inclusive_add_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_add_kernel_code_int, "test_wg_scan_inclusive_add_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_add_kernel_code_int,
+                                      "test_wg_scan_inclusive_add_int");
     if (err)
         return -1;
 
@@ -276,7 +278,9 @@ test_work_group_scan_inclusive_add_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_add_kernel_code_uint, "test_wg_scan_inclusive_add_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_add_kernel_code_uint,
+                                      "test_wg_scan_inclusive_add_uint");
     if (err)
         return -1;
 
@@ -378,7 +382,9 @@ test_work_group_scan_inclusive_add_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_add_kernel_code_long, "test_wg_scan_inclusive_add_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_add_kernel_code_long,
+                                      "test_wg_scan_inclusive_add_long");
     if (err)
         return -1;
 
@@ -481,7 +487,9 @@ test_work_group_scan_inclusive_add_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_add_kernel_code_ulong, "test_wg_scan_inclusive_add_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_add_kernel_code_ulong,
+                                      "test_wg_scan_inclusive_add_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_inclusive_max.cpp

@@ -175,7 +175,9 @@ test_work_group_scan_inclusive_max_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_max_kernel_code_int, "test_wg_scan_inclusive_max_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_max_kernel_code_int,
+                                      "test_wg_scan_inclusive_max_int");
     if (err)
         return -1;
 
@@ -278,7 +280,9 @@ test_work_group_scan_inclusive_max_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_max_kernel_code_uint, "test_wg_scan_inclusive_max_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_max_kernel_code_uint,
+                                      "test_wg_scan_inclusive_max_uint");
     if (err)
         return -1;
 
@@ -380,7 +384,9 @@ test_work_group_scan_inclusive_max_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_max_kernel_code_long, "test_wg_scan_inclusive_max_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_max_kernel_code_long,
+                                      "test_wg_scan_inclusive_max_long");
     if (err)
         return -1;
 
@@ -483,7 +489,9 @@ test_work_group_scan_inclusive_max_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_max_kernel_code_ulong, "test_wg_scan_inclusive_max_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_max_kernel_code_ulong,
+                                      "test_wg_scan_inclusive_max_ulong");
     if (err)
         return -1;
 

test_conformance/workgroups/test_wg_scan_inclusive_min.cpp

@@ -175,7 +175,9 @@ test_work_group_scan_inclusive_min_int(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_min_kernel_code_int, "test_wg_scan_inclusive_min_int", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_min_kernel_code_int,
+                                      "test_wg_scan_inclusive_min_int");
     if (err)
         return -1;
 
@@ -278,7 +280,9 @@ test_work_group_scan_inclusive_min_uint(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_min_kernel_code_uint, "test_wg_scan_inclusive_min_uint", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_min_kernel_code_uint,
+                                      "test_wg_scan_inclusive_min_uint");
     if (err)
         return -1;
 
@@ -380,7 +384,9 @@ test_work_group_scan_inclusive_min_long(cl_device_id device, cl_context context,
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_min_kernel_code_long, "test_wg_scan_inclusive_min_long", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_min_kernel_code_long,
+                                      "test_wg_scan_inclusive_min_long");
     if (err)
         return -1;
 
@@ -483,7 +489,9 @@ test_work_group_scan_inclusive_min_ulong(cl_device_id device, cl_context context
     int          i;
     MTdata       d;
 
-    err = create_single_kernel_helper_with_build_options( context, &program, &kernel, 1, &wg_scan_inclusive_min_kernel_code_ulong, "test_wg_scan_inclusive_min_ulong", "-cl-std=CL2.0" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &wg_scan_inclusive_min_kernel_code_ulong,
+                                      "test_wg_scan_inclusive_min_ulong");
     if (err)
         return -1;