Add ability to select data types, vector size and stride for basic tests

SlawomirMakowski · SlawomirMakowski · commit 3ff004edd87e · 2025-11-28T14:01:21.000+01:00
This commit introduces command-line options to the basic tests
that allow users to specify data types, vector sizes, and strides for the tests.
This enhancement provides flexibility and easier debugability in testing various configurations.

* added `v1|v2|v3|v4|v8|v16` options to select vector size
* added `s1|s3|s4|s5` options to select stride value
* added `char|uchar|short|ushort|int|uint|long|ulong|float|half|double` options to specify data type used in test.

Above options available for following tests:
- async_copy_global_to_local
- async_copy_local_to_global
- prefetch
- async_strided_copy_global_to_local
- async_strided_copy_local_to_global
diff --git a/test_conformance/basic/main.cpp b/test_conformance/basic/main.cpp
@@ -17,10 +17,16 @@
 #include "harness/deviceInfo.h"
 #include "harness/kernelHelpers.h"
 #include "harness/testHarness.h"
+#include "harness/parseParameters.h"
 
 #include <CL/cl_half.h>
+#include "testBase.h"
 
 cl_half_rounding_mode halfRoundingMode = CL_HALF_RTE;
+int gUseDataType = 0;
+unsigned int gUseVectorSize = 0;
+unsigned int gUseStride = 0;
+std::vector<ExplicitType> gVecType = {};
 
 test_status InitCL(cl_device_id device)
 {
@@ -46,8 +52,170 @@ test_status InitCL(cl_device_id device)
     return TEST_PASS;
 }
 
+static void PrintUsage(void)
+{
+    vlog("additional options for test cases:\n");
+    vlog("\t  async_copy_global_to_local\n\t  async_copy_local_to_global\n\t  "
+         "prefetch\n\t  async_strided_copy_global_to_local\n\t  "
+         "async_strided_copy_local_to_global\n\n");
+    vlog("\t   [v1|v2|v3|v4|v8|v16] - specify vector size used in test.\n");
+    vlog("\t   [char|uchar|short|ushort|int|uint|long|ulong|float|half|double] "
+         "- specify data type used in test.\n");
+    vlog("\n\t   In addition for async_strided_copy_global_to_local and "
+         "async_strided_copy_local_to_global:\n");
+    vlog("\t   [s1|s3|s4|s5] - specify stride size used in test.\n");
+    vlog("\n\t   Examples of usage:\n");
+    vlog("\t     test_basic async_copy_local_to_global float v4\n");
+    vlog("\t     test_basic prefetch int uint v2\n");
+    vlog("\t     test_basic async_strided_copy_global_to_local uchar s4 v1\n");
+    vlog("\t     test_basic async_strided_copy_local_to_global char s3 v4\n\n");
+    vlog("\t   Notes:\n");
+    vlog("\t     Choose exactly one vector size: v1|v2|v3|v4|v8|v16.\n");
+    vlog("\t     You can include multiple data types (e.g., char int).\n");
+    vlog("\t     For strided variants, include exactly one stride: "
+         "s1|s3|s4|s5.\n\n");
+    vlog("\n");
+}
+
 int main(int argc, const char *argv[])
 {
+    argc = parseCustomParam(argc, argv);
+
+    // Parse arguments
+    int argsRemoveNum = 0;
+    for (int i = 1; i < argc; i++)
+    {
+        if (strcmp(argv[i], "-h") == 0 || strcmp(argv[i], "--help") == 0)
+        {
+            PrintUsage();
+            return -1;
+        }
+        if (strcmp(argv[i], "char") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kChar);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "uchar") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kUChar);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "short") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kShort);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "ushort") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kUShort);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "int") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kInt);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "uint") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kUInt);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "long") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kLong);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "ulong") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kULong);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "float") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kFloat);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "half") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kHalf);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "double") == 0)
+        {
+            gUseDataType = 1;
+            gVecType.push_back(kDouble);
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v1") == 0)
+        {
+            gUseVectorSize = 1;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v2") == 0)
+        {
+            gUseVectorSize = 2;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v3") == 0)
+        {
+            gUseVectorSize = 3;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v4") == 0)
+        {
+            gUseVectorSize = 4;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v8") == 0)
+        {
+            gUseVectorSize = 8;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "v16") == 0)
+        {
+            gUseVectorSize = 16;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "s1") == 0)
+        {
+            gUseStride = 1;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "s3") == 0)
+        {
+            gUseStride = 3;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "s4") == 0)
+        {
+            gUseStride = 4;
+            argsRemoveNum += 1;
+        }
+        else if (strcmp(argv[i], "s5") == 0)
+        {
+            gUseStride = 5;
+            argsRemoveNum += 1;
+        }
+    }
+        // remove additionally parsed args from argv
+        if (argsRemoveNum > 0)
+        {
+        for (int j = argc; j < argc - argsRemoveNum; j++)
+        {
+                argv[j] = argv[j + argsRemoveNum];
+            }
+            argc -= argsRemoveNum;
+        }
     return runTestHarnessWithCheck(
         argc, argv, test_registry::getInstance().num_tests(),
         test_registry::getInstance().definitions(), false, 0, InitCL);
diff --git a/test_conformance/basic/testBase.h b/test_conformance/basic/testBase.h
@@ -22,5 +22,9 @@
 #include "harness/testHarness.h"
 #include "harness/typeWrappers.h"
 #include "harness/rounding_mode.h"
+extern int gUseDataType;
+extern unsigned int gUseVectorSize;
+extern unsigned int gUseStride;
+extern std::vector<ExplicitType> gVecType;
 
 #endif // _testBase_h
diff --git a/test_conformance/basic/test_async_copy.cpp b/test_conformance/basic/test_async_copy.cpp
@@ -23,51 +23,51 @@
 #include "testBase.h"
 
 static const char *async_global_to_local_kernel =
-"%s\n" // optional pragma string
+    "%s\n" // optional pragma string
 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
-"{\n"
-" int i;\n"
-// Zero the local storage first
-" for(i=0; i<copiesPerWorkItem; i++)\n"
-"     localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
+    "{\n"
+    " int i;\n"
+    // Zero the local storage first
+    " for(i=0; i<copiesPerWorkItem; i++)\n"
+    "     localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
 // Do this to verify all kernels are done zeroing the local buffer before we try the copy
-"    barrier( CLK_LOCAL_MEM_FENCE );\n"
-"    event_t event;\n"
+    "    barrier( CLK_LOCAL_MEM_FENCE );\n"
+    "    event_t event;\n"
 "    event = async_work_group_copy( (__local %s*)localBuffer, (__global const %s*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, 0 );\n"
 // Wait for the copy to complete, then verify by manually copying to the dest
-"    wait_group_events( 1, &event );\n"
-" for(i=0; i<copiesPerWorkItem; i++)\n"
+    "    wait_group_events( 1, &event );\n"
+    " for(i=0; i<copiesPerWorkItem; i++)\n"
 "  dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ];\n"
 "}\n" ;
 
 static const char *async_local_to_global_kernel =
-"%s\n" // optional pragma string
+    "%s\n" // optional pragma string
 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
-"{\n"
-" int i;\n"
-// Zero the local storage first
-" for(i=0; i<copiesPerWorkItem; i++)\n"
-"  localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
+    "{\n"
+    " int i;\n"
+    // Zero the local storage first
+    " for(i=0; i<copiesPerWorkItem; i++)\n"
+    "  localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
 // Do this to verify all kernels are done zeroing the local buffer before we try the copy
-"    barrier( CLK_LOCAL_MEM_FENCE );\n"
-" for(i=0; i<copiesPerWorkItem; i++)\n"
+    "    barrier( CLK_LOCAL_MEM_FENCE );\n"
+    " for(i=0; i<copiesPerWorkItem; i++)\n"
 "  localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = src[ get_global_id( 0 )*copiesPerWorkItem+i ];\n"
 // Do this to verify all kernels are done copying to the local buffer before we try the copy
-"    barrier( CLK_LOCAL_MEM_FENCE );\n"
-"    event_t event;\n"
+    "    barrier( CLK_LOCAL_MEM_FENCE );\n"
+    "    event_t event;\n"
 "    event = async_work_group_copy((__global %s*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const %s*)localBuffer, (size_t)copiesPerWorkgroup, 0 );\n"
-"    wait_group_events( 1, &event );\n"
+    "    wait_group_events( 1, &event );\n"
 "}\n" ;
 
 
 static const char *prefetch_kernel =
-"%s\n" // optional pragma string
+    "%s\n" // optional pragma string
 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
-"{\n"
-" // Ignore this: %s%s%s\n"
-" int i;\n"
+    "{\n"
+    " // Ignore this: %s%s%s\n"
+    " int i;\n"
 " prefetch( (const __global %s*)(src+copiesPerWorkItem*get_global_id(0)), copiesPerWorkItem);\n"
-" for(i=0; i<copiesPerWorkItem; i++)\n"
+    " for(i=0; i<copiesPerWorkItem; i++)\n"
 "  dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = src[ get_global_id( 0 )*copiesPerWorkItem+i ];\n"
 "}\n" ;
 
@@ -237,12 +237,30 @@ static int test_copy(cl_device_id deviceID, cl_context context,
 static int test_copy_all_types(cl_device_id deviceID, cl_context context,
                                cl_command_queue queue, const char *kernelCode)
 {
-    const std::vector<ExplicitType> vecType = { kChar,  kUChar, kShort, kUShort,
-                                                kInt,   kUInt,  kLong,  kULong,
-                                                kFloat, kHalf,  kDouble };
-    unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
+    std::vector<ExplicitType> vecType = { kChar,  kUChar, kShort, kUShort,
+                                          kInt,   kUInt,  kLong,  kULong,
+                                          kFloat, kHalf,  kDouble };
     unsigned int size, typeIndex;
 
+    std::vector<unsigned int> vecSizes;
+
+    if (gUseDataType)
+    {
+        log_info("WARNING: Running subset of data types !\n");
+        vecType = gVecType;
+    }
+
+    if (gUseVectorSize)
+    {
+        log_info("WARNING: Running for subset of vector size: %d only !\n",
+                 gUseVectorSize);
+        vecSizes.assign({ gUseVectorSize, 0 });
+    }
+    else
+    {
+        vecSizes.assign({ 1, 2, 3, 4, 8, 16, 0 });
+    }
+
     int errors = 0;
 
     bool fp16Support = is_extension_available(deviceID, "cl_khr_fp16");
diff --git a/test_conformance/basic/test_async_strided_copy.cpp b/test_conformance/basic/test_async_strided_copy.cpp
@@ -237,12 +237,38 @@ static int test_strided_copy_all_types(cl_device_id deviceID,
                                        cl_command_queue queue,
                                        const char *kernelCode)
 {
-    const std::vector<ExplicitType> vecType = { kChar,  kUChar, kShort, kUShort,
+    std::vector<ExplicitType> vecType = { kChar,  kUChar, kShort, kUShort,
                                                 kInt,   kUInt,  kLong,  kULong,
                                                 kFloat, kHalf,  kDouble };
-    const unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
-    const unsigned int strideSizes[] = { 1, 3, 4, 5, 0 };
     unsigned int size, typeIndex, stride;
+    std::vector<unsigned int> vecSizes;
+    std::vector<unsigned int> strideSizes;
+
+    if (gUseDataType)
+    {
+        log_info("WARNING: Running subset of data types !\n");
+        vecType = gVecType;
+    }
+
+    if (gUseVectorSize)
+    {
+        log_info("WARNING: Running for vector size: %d only !\n",
+                 gUseVectorSize);
+        vecSizes.assign({ gUseVectorSize, 0 });
+    }
+    else
+    {
+        vecSizes.assign({ 1, 2, 3, 4, 8, 16, 0 });
+    }
+    if (gUseStride)
+    {
+        log_info("WARNING: Running for stride: %d only !\n", gUseStride);
+        strideSizes.assign({ gUseStride, 0 });
+    }
+    else
+    {
+        strideSizes.assign({ 1, 3, 4, 5, 0 });
+    }
 
     int errors = 0;
 
