Merge branch 'main' into mutable_command_cross_queue_simultaneous_use

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h
shajder · Mar 15, 2024 · 7eaff77 · 7eaff77
2 parents d080e3e + d8ad10d
commit 7eaff77
Show file tree

Hide file tree

Showing 20 changed files with 586 additions and 46 deletions.
diff --git a/test_conformance/api/test_mem_object_info.cpp b/test_conformance/api/test_mem_object_info.cpp
@@ -223,6 +223,12 @@ int test_get_buffer_info( cl_device_id deviceID, cl_context context, cl_command_
         }
 
         // Perform buffer object queries.
+        void *ptr;
+        TEST_MEM_OBJECT_PARAM(
+            bufferObject, CL_MEM_HOST_PTR, ptr,
+            ((bufferFlags[i] & CL_MEM_USE_HOST_PTR) ? buffer : NULL),
+            "host pointer", "%p", void *)
+
         cl_mem_object_type type;
         TEST_MEM_OBJECT_PARAM( bufferObject, CL_MEM_TYPE, type, CL_MEM_OBJECT_BUFFER, "type", "%d", int )
 

diff --git a/test_conformance/basic/test_async_copy_fence.cpp b/test_conformance/basic/test_async_copy_fence.cpp
@@ -214,15 +214,15 @@ static const char *export_after_import_aliased_local_kernel =
     "    event_t events;\n"
     "    events = async_work_group_copy( (__local %s*)localBuffer, (__global "
     "const %s*)(importSrc+importSrcLocalSize*get_group_id(0)), "
-    "(size_t)importSrcLocalSize, events );\n"
+    "(size_t)importSrcLocalSize, 0 );\n"
     "    async_work_group_copy_fence( CLK_LOCAL_MEM_FENCE );\n"
     "    events = async_work_group_copy((__global "
     "%s*)(exportDst+exportSrcLocalSize*get_group_id(0)), (__local const "
     "%s*)(localBuffer + (importSrcLocalSize - exportSrcLocalSize)), "
-    "(size_t)exportSrcLocalSize, 0 );\n"
+    "(size_t)exportSrcLocalSize, events );\n"
     // Wait for the import and export to complete, then verify by manually
     // copying to the dest
-    "    wait_group_events( 2, &events );\n"
+    "    wait_group_events( 1, &events );\n"
     "    for(i=0; i<importCopiesPerWorkItem; i++) {\n"
     "        importDst[ get_global_id( 0 )*importCopiesPerWorkItem+i ] = "
     "localBuffer[ get_local_id( 0 )*importCopiesPerWorkItem+i ];\n"

diff --git a/test_conformance/device_execution/device_queue.cpp b/test_conformance/device_execution/device_queue.cpp
@@ -84,7 +84,7 @@ static int check_device_queues(cl_device_id device, cl_context context, cl_uint
     for(i = 0; i < num_queues; ++i)
     {
         queue[i] = clCreateCommandQueueWithProperties(context, device, properties, &err_ret);
-        test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE) failed");
+        test_error(err_ret, "clCreateCommandQueueWithProperties failed");
     }
 
     // Validate all queues
@@ -128,7 +128,9 @@ int test_device_queue(cl_device_id device, cl_context context, cl_command_queue
     if(max_queues > MAX_QUEUES) max_queues = MAX_QUEUES;
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     err_ret = check_device_queue(device, context, dev_queue, preffered_size);
     if(check_error(err_ret, "Default device queue validation failed")) res = -1;

diff --git a/test_conformance/device_execution/enqueue_block.cpp b/test_conformance/device_execution/enqueue_block.cpp
@@ -623,7 +623,9 @@ int test_enqueue_block(cl_device_id device, cl_context context, cl_command_queue
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     size_t global_size = MAX_GWS;
     size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;

diff --git a/test_conformance/device_execution/enqueue_flags.cpp b/test_conformance/device_execution/enqueue_flags.cpp
@@ -696,7 +696,9 @@ int test_enqueue_flags(cl_device_id device, cl_context context, cl_command_queue
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     size_t global_size = MAX_GWS;
     size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;

diff --git a/test_conformance/device_execution/enqueue_multi_queue.cpp b/test_conformance/device_execution/enqueue_multi_queue.cpp
@@ -123,7 +123,9 @@ int test_enqueue_multi_queue(cl_device_id device, cl_context context, cl_command
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     if(max_queues > 1)
     {
@@ -141,7 +143,9 @@ int test_enqueue_multi_queue(cl_device_id device, cl_context context, cl_command
         for(i = 0; i < n; ++i)
         {
             queues[i] = clCreateCommandQueueWithProperties(context, device, queue_prop, &err_ret);
-            test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE) failed");
+            test_error(err_ret,
+                       "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE) "
+                       "failed");
             q[i] = queues[i];
         }
 

diff --git a/test_conformance/device_execution/enqueue_ndrange.cpp b/test_conformance/device_execution/enqueue_ndrange.cpp
@@ -641,7 +641,9 @@ int test_enqueue_ndrange(cl_device_id device, cl_context context, cl_command_que
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     max_local_size = (max_local_size > MAX_GWS)? MAX_GWS: max_local_size;
     if(gWimpyMode)

diff --git a/test_conformance/device_execution/enqueue_profiling.cpp b/test_conformance/device_execution/enqueue_profiling.cpp
@@ -66,18 +66,18 @@ int test_enqueue_profiling(cl_device_id device, cl_context context,
     dev_queue = clCreateCommandQueueWithProperties(
         context, device, dev_queue_prop_def, &err_ret);
     test_error(err_ret,
-               "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_"
-               "DEFAULT) failed");
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     cl_queue_properties host_queue_prop_def[] = { CL_QUEUE_PROPERTIES,
                                                   CL_QUEUE_PROFILING_ENABLE,
                                                   0 };
 
     host_queue = clCreateCommandQueueWithProperties(
         context, device, host_queue_prop_def, &err_ret);
-    test_error(err_ret,
-               "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_"
-               "DEFAULT) failed");
+    test_error(
+        err_ret,
+        "clCreateCommandQueueWithProperties(CL_QUEUE_PROFILING_ENABLE) failed");
 
     cl_int status;
     size_t size = 1;

diff --git a/test_conformance/device_execution/enqueue_wg_size.cpp b/test_conformance/device_execution/enqueue_wg_size.cpp
@@ -1659,7 +1659,9 @@ int test_enqueue_wg_size(cl_device_id device, cl_context context, cl_command_que
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
 
     size_t failCnt = 0;

diff --git a/test_conformance/device_execution/host_multi_queue.cpp b/test_conformance/device_execution/host_multi_queue.cpp
@@ -157,7 +157,9 @@ int test_host_multi_queue(cl_device_id device, cl_context context, cl_command_qu
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     cl_uint n = num_kernels_multi_queue_block; // Number of host queues
     std::vector<clCommandQueueWrapper> queues(n);

diff --git a/test_conformance/device_execution/host_queue_order.cpp b/test_conformance/device_execution/host_queue_order.cpp
@@ -117,7 +117,9 @@ int test_host_queue_order(cl_device_id device, cl_context context, cl_command_qu
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     cl_int status;
     size_t size = 1;

diff --git a/test_conformance/device_execution/nested_blocks.cpp b/test_conformance/device_execution/nested_blocks.cpp
@@ -335,7 +335,9 @@ int test_enqueue_nested_blocks(cl_device_id device, cl_context context, cl_comma
     };
 
     dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
-    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");
+    test_error(err_ret,
+               "clCreateCommandQueueWithProperties(CL_QUEUE_ON_DEVICE | "
+               "CL_QUEUE_ON_DEVICE_DEFAULT) failed");
 
     kernel_arg args[] =
     {

diff --git a/...ce/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt b/...ce/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt
@@ -9,6 +9,7 @@ set(${MODULE_NAME}_SOURCES
     mutable_command_global_size.cpp
     mutable_command_local_size.cpp
     mutable_command_global_offset.cpp
+    mutable_command_full_dispatch.cpp
     ../basic_command_buffer.cpp
 )
 

diff --git a/...formance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp b/...formance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp
@@ -26,6 +26,7 @@ test_definition test_list[] = {
     ADD_TEST(mutable_command_info_global_work_offset),
     ADD_TEST(mutable_command_info_local_work_size),
     ADD_TEST(mutable_command_info_global_work_size),
+    ADD_TEST(mutable_command_full_dispatch),
     ADD_TEST(mutable_dispatch_image_1d_arguments),
     ADD_TEST(mutable_dispatch_image_2d_arguments),
     ADD_TEST(mutable_dispatch_out_of_order),