Merge branch 'main' into pr/execute-multipass

KhronosGroup · Feb 10, 2025 · ef84ed6 · ef84ed6
2 parents 3d0cd94 + a8b32b2
commit ef84ed6
Show file tree

Hide file tree

Showing 112 changed files with 1,680 additions and 1,159 deletions.
diff --git a/.gitignore b/.gitignore
@@ -0,0 +1,2 @@
+# Build dir
+[Bb]uild/
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -107,6 +107,7 @@ if(CMAKE_COMPILER_IS_GNUCC OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "(Apple)?Clang"
     add_cxx_flag_if_supported(-Wno-error=cpp) # Allow #warning directive
     add_cxx_flag_if_supported(-Wno-unknown-pragmas) # Issue #785
     add_cxx_flag_if_supported(-Wno-error=asm-operand-widths) # Issue #784
+    add_cxx_flag_if_supported(-Wno-strict-aliasing)  # Issue 2234
 
     # -msse -mfpmath=sse to force gcc to use sse for float math,
     # avoiding excess precision problems that cause tests like int2float

diff --git a/test_common/harness/errorHelpers.cpp b/test_common/harness/errorHelpers.cpp
@@ -111,6 +111,7 @@ const char *IGetErrorString(int clErrorCode)
             return "CL_INVALID_SYNC_POINT_WAIT_LIST_KHR";
         case CL_INVALID_COMMAND_BUFFER_KHR:
             return "CL_INVALID_COMMAND_BUFFER_KHR";
+        case CL_INVALID_SEMAPHORE_KHR: return "CL_INVALID_SEMAPHORE_KHR";
         default: return "(unknown)";
     }
 }

diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp
@@ -2415,6 +2415,12 @@ int debug_find_vector_in_image(void *imagePtr, image_descriptor *imageInfo,
                 (imageInfo->height >> lod) ? (imageInfo->height >> lod) : 1;
             depth = (imageInfo->depth >> lod) ? (imageInfo->depth >> lod) : 1;
             break;
+        default:
+            log_error("ERROR: Invalid imageInfo->type = %d\n", imageInfo->type);
+            width = 0;
+            depth = 0;
+            height = 0;
+            break;
     }
 
     row_pitch = width * get_pixel_size(imageInfo->format);
@@ -3661,6 +3667,11 @@ void copy_image_data(image_descriptor *srcImageInfo,
                     ? (srcImageInfo->height >> src_lod)
                     : 1;
                 break;
+            default:
+                log_error("ERROR: Invalid srcImageInfo->type = %d\n",
+                          srcImageInfo->type);
+                src_lod = 0;
+                break;
         }
         src_mip_level_offset = compute_mip_level_offset(srcImageInfo, src_lod);
         src_row_pitch_lod =
@@ -3707,6 +3718,11 @@ void copy_image_data(image_descriptor *srcImageInfo,
                     ? (dstImageInfo->height >> dst_lod)
                     : 1;
                 break;
+            default:
+                log_error("ERROR: Invalid dstImageInfo->num_mip_levels = %d\n",
+                          dstImageInfo->num_mip_levels);
+                dst_lod = 0;
+                break;
         }
         dst_mip_level_offset = compute_mip_level_offset(dstImageInfo, dst_lod);
         dst_row_pitch_lod =

diff --git a/test_common/harness/typeWrappers.cpp b/test_common/harness/typeWrappers.cpp
@@ -348,7 +348,7 @@ cl_int clProtectedImage::Create(cl_context context,
                                 const cl_image_format *fmt, size_t width,
                                 size_t height, size_t depth, size_t arraySize)
 {
-    cl_int error;
+    cl_int error = 0;
 #if defined(__APPLE__)
     int protect_pages = 1;
     cl_device_id devices[16];

diff --git a/test_conformance/SVM/test_enqueue_api.cpp b/test_conformance/SVM/test_enqueue_api.cpp
@@ -160,8 +160,8 @@ REGISTER_TEST(svm_enqueue_api)
             error = clSetUserEventStatus(userEvent, CL_COMPLETE);
             test_error(error, "clSetUserEventStatus failed");
 
-            cl_uchar *src_ptr;
-            cl_uchar *dst_ptr;
+            cl_uchar *src_ptr = nullptr;
+            cl_uchar *dst_ptr = nullptr;
             if (test_case.srcAlloc == host)
             {
                 src_ptr = srcHostData.data();

diff --git a/test_conformance/api/test_queue.cpp b/test_conformance/api/test_queue.cpp
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2020 The Khronos Group Inc.
+// Copyright (c) 2024 The Khronos Group Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -58,3 +58,61 @@ REGISTER_TEST(queue_flush_on_release)
 
     return success ? TEST_PASS : TEST_FAIL;
 }
+
+REGISTER_TEST(multi_queue_flush_on_release)
+{
+    cl_int err;
+
+    // Create A command queue
+    cl_command_queue queue_A = clCreateCommandQueue(context, device, 0, &err);
+    test_error(err, "Could not create command queue A");
+
+    // Create B command queue
+    cl_command_queue queue_B = clCreateCommandQueue(context, device, 0, &err);
+    test_error(err, "Could not create command queue B");
+
+    // Create a kernel
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+    const char *source = "void kernel test(){}";
+    err = create_single_kernel_helper(context, &program, &kernel, 1, &source,
+                                      "test");
+    test_error(err, "Could not create kernel");
+
+    // Enqueue the kernel on queue_A and obtain event to synchronize with
+    // queue_B
+    size_t gws = num_elements;
+    clEventWrapper event_A;
+    err = clEnqueueNDRangeKernel(queue_A, kernel, 1, nullptr, &gws, nullptr, 0,
+                                 nullptr, &event_A);
+    test_error(err, "Could not enqueue kernel");
+
+    // Enqueue the kernel on queue_B using event_A for synchronization and
+    // create event_B to track completion
+    clEventWrapper event_B;
+    err = clEnqueueNDRangeKernel(queue_B, kernel, 1, nullptr, &gws, nullptr, 1,
+                                 &event_A, &event_B);
+    test_error(err, "Could not enqueue kernel");
+
+    // Release queue_A, which performs an implicit flush to issue any previously
+    // queued OpenCL commands
+    err = clReleaseCommandQueue(queue_A);
+    test_error(err, "clReleaseCommandQueue failed");
+
+    err = clFlush(queue_B);
+    test_error(err, "clFlush failed");
+
+    // Wait for kernel to execute since the queue must flush on release
+    bool success = poll_until(2000, 50, [&event_B]() {
+        cl_int status;
+        cl_int err = clGetEventInfo(event_B, CL_EVENT_COMMAND_EXECUTION_STATUS,
+                                    sizeof(cl_int), &status, nullptr);
+        if ((err != CL_SUCCESS) || (status != CL_COMPLETE))
+        {
+            return false;
+        }
+        return true;
+    });
+
+    return success ? TEST_PASS : TEST_FAIL;
+}
diff --git a/test_conformance/api/test_wg_suggested_local_work_size.cpp b/test_conformance/api/test_wg_suggested_local_work_size.cpp
@@ -208,7 +208,7 @@ int do_test_work_group_suggested_local_size(
     bool (*skip_cond)(size_t), size_t start, size_t end, size_t incr,
     cl_ulong max_local_mem_size, size_t global_work_offset[], num_dims dim)
 {
-    int err;
+    int err = 0;
     size_t test_values[] = { 1, 1, 1 };
     std::string kernel_names[6] = {
         "test_wg_scan_local_work_group_size",

diff --git a/test_conformance/basic/main.cpp b/test_conformance/basic/main.cpp
@@ -99,6 +99,8 @@ test_definition test_list[] = {
     ADD_TEST(enqueue_map_image),
 
     ADD_TEST(work_item_functions),
+    ADD_TEST(work_item_functions_out_of_range),
+    ADD_TEST(work_item_functions_out_of_range_hardcoded),
 
     ADD_TEST(astype),
 

diff --git a/test_conformance/basic/procs.h b/test_conformance/basic/procs.h
@@ -98,6 +98,13 @@ extern int      test_enqueue_map_buffer(cl_device_id deviceID, cl_context contex
 extern int      test_enqueue_map_image(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
 
 extern int      test_work_item_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
+extern int test_work_item_functions_out_of_range(cl_device_id deviceID,
+                                                 cl_context context,
+                                                 cl_command_queue queue,
+                                                 int num_elements);
+extern int test_work_item_functions_out_of_range_hardcoded(
+    cl_device_id deviceID, cl_context context, cl_command_queue queue,
+    int num_elements);
 
 extern int      test_astype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
 

diff --git a/test_conformance/basic/test_imagereadwrite.cpp b/test_conformance/basic/test_imagereadwrite.cpp
@@ -314,6 +314,12 @@ test_imagereadwrite(cl_device_id device, cl_context context, cl_command_queue qu
                 }
                 outp = (void *)rgbafp_outptr;
                 break;
+            default:
+                log_error("ERROR Invalid j = %d\n", j);
+                elem_size = 0;
+                p = nullptr;
+                outp = nullptr;
+                break;
         }
 
         const char* update_packed_pitch_name = "";

diff --git a/test_conformance/basic/test_imagereadwrite3d.cpp b/test_conformance/basic/test_imagereadwrite3d.cpp
@@ -320,6 +320,12 @@ test_imagereadwrite3d(cl_device_id device, cl_context context, cl_command_queue
                 }
                 outp = (void *)rgbafp_outptr;
                 break;
+            default:
+                log_error("ERROR Invalid j = %d\n", j);
+                elem_size = 0;
+                p = nullptr;
+                outp = nullptr;
+                break;
         }
 
         const char* update_packed_pitch_name = "";