Merge branch 'main' into command_buffer_req_group_size

Author: bashbaug

test_common/harness/deviceInfo.cpp

@@ -50,7 +50,7 @@ std::string get_device_info_string(cl_device_id device,
 }
 
 /* Determines if an extension is supported by a device. */
-int is_extension_available(cl_device_id device, const char *extensionName)
+bool is_extension_available(cl_device_id device, const char *extensionName)
 {
     std::string extString = get_device_extensions_string(device);
     std::istringstream ss(extString);

test_common/harness/deviceInfo.h

@@ -26,7 +26,7 @@ std::string get_device_info_string(cl_device_id device,
                                    cl_device_info param_name);
 
 /* Determines if an extension is supported by a device. */
-int is_extension_available(cl_device_id device, const char *extensionName);
+bool is_extension_available(cl_device_id device, const char *extensionName);
 
 /* Returns the version of the extension the device supports or throws an
  * exception if the extension is not supported by the device. */

test_conformance/api/CMakeLists.txt

@@ -6,6 +6,7 @@ set(${MODULE_NAME}_SOURCES
          main.cpp
          negative_platform.cpp
          negative_queue.cpp
+         negative_enqueue_marker.cpp
          negative_enqueue_map_image.cpp
          test_api_consistency.cpp
          test_bool.cpp

test_conformance/api/negative_enqueue_marker.cpp

@@ -0,0 +1,97 @@
+//
+// Copyright (c) 2025 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.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "testBase.h"
+#include "harness/typeWrappers.h"
+
+REGISTER_TEST(negative_enqueue_marker_with_wait_list)
+{
+    cl_platform_id platform = getPlatformFromDevice(device);
+    cl_context_properties props[3] = {
+        CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platform),
+        0
+    };
+
+    cl_int err = CL_SUCCESS;
+    clContextWrapper ctx =
+        clCreateContext(props, 1, &device, nullptr, nullptr, &err);
+    test_error(err, "clCreateContext failed");
+
+    cl_event ret_event = nullptr;
+
+    err = clEnqueueMarkerWithWaitList(nullptr, 0, nullptr, &ret_event);
+    test_failure_error_ret(err, CL_INVALID_COMMAND_QUEUE,
+                           "clEnqueueMarkerWithWaitList should return "
+                           "CL_INVALID_COMMAND_QUEUE when: \"command_queue is "
+                           "not a valid host command-queue\" using a nullptr",
+                           TEST_FAIL);
+    test_assert_error(ret_event == nullptr,
+                      "if clEnqueueMarkerWithWaitList failed, no ret_event "
+                      "should be created");
+
+    clEventWrapper different_ctx_event = clCreateUserEvent(ctx, &err);
+    test_error(err, "clCreateUserEvent failed");
+
+    err =
+        clEnqueueMarkerWithWaitList(queue, 1, &different_ctx_event, &ret_event);
+    test_failure_error_ret(
+        err, CL_INVALID_CONTEXT,
+        "clEnqueueMarkerWithWaitList should return CL_INVALID_CONTEXT when: "
+        "\"The context of both the command queue and the events in ret_event "
+        "wait list are not the same\"",
+        TEST_FAIL);
+    test_assert_error(ret_event == nullptr,
+                      "if clEnqueueMarkerWithWaitList failed, no ret_event "
+                      "should be created");
+
+    err = clEnqueueMarkerWithWaitList(queue, 1, nullptr, &ret_event);
+    test_failure_error_ret(
+        err, CL_INVALID_EVENT_WAIT_LIST,
+        "clEnqueueMarkerWithWaitList should return CL_INVALID_EVENT_WAIT_LIST "
+        "when: \"num_events_in_wait_list > 0 but event_wait_list is NULL\"",
+        TEST_FAIL);
+    test_assert_error(ret_event == nullptr,
+                      "if clEnqueueMarkerWithWaitList failed, no ret_event "
+                      "should be created");
+
+
+    clEventWrapper event = clCreateUserEvent(context, &err);
+    test_error(err, "clCreateUserEvent failed");
+
+    err = clEnqueueMarkerWithWaitList(queue, 0, &event, &ret_event);
+    test_failure_error_ret(
+        err, CL_INVALID_EVENT_WAIT_LIST,
+        "clEnqueueMarkerWithWaitList should return CL_INVALID_EVENT_WAIT_LIST "
+        "when: \"num_events_in_wait_list is 0 but event_wait_list is not "
+        "NULL\"",
+        TEST_FAIL);
+    test_assert_error(ret_event == nullptr,
+                      "if clEnqueueMarkerWithWaitList failed, no ret_event "
+                      "should be created");
+
+    cl_event invalid_event_wait_list[] = { nullptr };
+    err = clEnqueueMarkerWithWaitList(queue, 1, invalid_event_wait_list,
+                                      &ret_event);
+    test_failure_error_ret(
+        err, CL_INVALID_EVENT_WAIT_LIST,
+        "clEnqueueMarkerWithWaitList should return CL_INVALID_EVENT_WAIT_LIST "
+        "when: \"event objects in event_wait_list are not valid events\"",
+        TEST_FAIL);
+    test_assert_error(ret_event == nullptr,
+                      "if clEnqueueMarkerWithWaitList failed, no ret_event "
+                      "should be created");
+
+    return TEST_PASS;
+}

test_conformance/api/test_kernel_attributes.cpp

@@ -386,14 +386,14 @@ REGISTER_TEST(null_required_work_group_size)
 
     struct KernelAttribInfo
     {
-        std::string str;
-        cl_uint max_dim;
+        cl_int wgs[3];
+        cl_uint min_dim;
     };
 
     std::vector<KernelAttribInfo> attribs;
-    attribs.push_back({ "__attribute__((reqd_work_group_size(2,1,1)))", 1 });
-    attribs.push_back({ "__attribute__((reqd_work_group_size(2,3,1)))", 2 });
-    attribs.push_back({ "__attribute__((reqd_work_group_size(2,3,4)))", 3 });
+    attribs.push_back({ { 2, 1, 1 }, 1 });
+    attribs.push_back({ { 2, 3, 1 }, 2 });
+    attribs.push_back({ { 2, 3, 4 }, 3 });
 
     const std::string body_str = R"(
         __kernel void wg_size(__global int* dst)
@@ -410,7 +410,11 @@ REGISTER_TEST(null_required_work_group_size)
 
     for (auto& attrib : attribs)
     {
-        const std::string source_str = attrib.str + body_str;
+        const std::string attrib_str = "__attribute__((reqd_work_group_size("
+            + std::to_string(attrib.wgs[0]) + ","
+            + std::to_string(attrib.wgs[1]) + ","
+            + std::to_string(attrib.wgs[2]) + ")))";
+        const std::string source_str = attrib_str + body_str;
         const char* source = source_str.c_str();
 
         clProgramWrapper program;
@@ -422,21 +426,19 @@ REGISTER_TEST(null_required_work_group_size)
         error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dst);
         test_error(error, "clSetKernelArg failed");
 
-        for (cl_uint work_dim = 1; work_dim <= attrib.max_dim; work_dim++)
+        for (cl_uint work_dim = attrib.min_dim; work_dim <= 3; work_dim++)
         {
-            const cl_int expected[3] = { 2, work_dim >= 2 ? 3 : 1,
-                                         work_dim >= 3 ? 4 : 1 };
             const size_t test_work_group_size =
-                expected[0] * expected[1] * expected[2];
-            if ((size_t)expected[0] > device_max_work_item_sizes[0]
-                || (size_t)expected[1] > device_max_work_item_sizes[1]
-                || (size_t)expected[2] > device_max_work_item_sizes[2]
+                attrib.wgs[0] * attrib.wgs[1] * attrib.wgs[2];
+            if ((size_t)attrib.wgs[0] > device_max_work_item_sizes[0]
+                || (size_t)attrib.wgs[1] > device_max_work_item_sizes[1]
+                || (size_t)attrib.wgs[2] > device_max_work_item_sizes[2]
                 || test_work_group_size > device_max_work_group_size)
             {
                 log_info("Skipping test for work_dim = %u: required work group "
                          "size (%i, %i, %i) (total %zu) exceeds device max "
                          "work group size (%zu, %zu, %zu) (total %zu)\n",
-                         work_dim, expected[0], expected[1], expected[2],
+                         work_dim, attrib.wgs[0], attrib.wgs[1], attrib.wgs[2],
                          test_work_group_size, device_max_work_item_sizes[0],
                          device_max_work_item_sizes[1],
                          device_max_work_item_sizes[2],
@@ -445,8 +447,9 @@ REGISTER_TEST(null_required_work_group_size)
             }
 
             const cl_int zero = 0;
-            error = clEnqueueFillBuffer(queue, dst, &zero, sizeof(zero), 0,
-                                        sizeof(expected), 0, nullptr, nullptr);
+            error =
+                clEnqueueFillBuffer(queue, dst, &zero, sizeof(zero), 0,
+                                    sizeof(attrib.wgs), 0, nullptr, nullptr);
             test_error(error, "clEnqueueFillBuffer failed");
 
             const size_t global_work_size[3] = { 2 * 32, 3 * 32, 4 * 32 };
@@ -460,12 +463,12 @@ REGISTER_TEST(null_required_work_group_size)
                                         results, 0, nullptr, nullptr);
             test_error(error, "clEnqueueReadBuffer failed");
 
-            if (results[0] != expected[0] || results[1] != expected[1]
-                || results[2] != expected[2])
+            if (results[0] != attrib.wgs[0] || results[1] != attrib.wgs[1]
+                || results[2] != attrib.wgs[2])
             {
                 log_error("Executed local size mismatch with work_dim = %u: "
                           "Expected (%d,%d,%d) got (%d,%d,%d)\n",
-                          work_dim, expected[0], expected[1], expected[2],
+                          work_dim, attrib.wgs[0], attrib.wgs[1], attrib.wgs[2],
                           results[0], results[1], results[2]);
                 return TEST_FAIL;
             }
@@ -479,15 +482,15 @@ REGISTER_TEST(null_required_work_group_size)
                 test_error(error,
                            "clGetKernelSuggestedLocalWorkSizeKHR failed");
 
-                if ((cl_int)suggested[0] != expected[0]
-                    || (cl_int)suggested[1] != expected[1]
-                    || (cl_int)suggested[2] != expected[2])
+                if (suggested[0] != (size_t)attrib.wgs[0]
+                    || suggested[1] != (size_t)attrib.wgs[1]
+                    || suggested[2] != (size_t)attrib.wgs[2])
                 {
                     log_error("Suggested local size mismatch with work_dim = "
-                              "%u: Expected (%d,%d,%d) got (%d,%d,%d)\n",
-                              work_dim, expected[0], expected[1], expected[2],
-                              (cl_int)suggested[0], (cl_int)suggested[1],
-                              (cl_int)suggested[2]);
+                              "%u: Expected (%d,%d,%d) got (%zu,%zu,%zu)\n",
+                              work_dim, attrib.wgs[0], attrib.wgs[1],
+                              attrib.wgs[2], suggested[0], suggested[1],
+                              suggested[2]);
                     return TEST_FAIL;
                 }
             }

test_conformance/api/test_kernels.cpp

@@ -982,6 +982,12 @@ REGISTER_TEST(negative_invalid_arg_index)
 
 REGISTER_TEST(negative_invalid_arg_size_local)
 {
+    if (true)
+    {
+        log_info("Disabling this test temporarily, see internal issue 374.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
     cl_int error = CL_SUCCESS;
     clProgramWrapper program;
     clKernelWrapper local_arg_kernel;

test_conformance/c11_atomics/common.h

@@ -702,6 +702,28 @@ class CBasicTestMemOrder2Scope
                                            cl_context context,
                                            cl_command_queue queue)
     {
+        // Comparator for orders and scopes.
+        const auto checkValidity = [](TExplicitMemoryOrderType success,
+                                      TExplicitMemoryOrderType failure,
+                                      TExplicitMemoryScopeType scope) {
+            // Both memory order arguments must be set (or neither).
+            if ((success == MEMORY_ORDER_EMPTY || failure == MEMORY_ORDER_EMPTY)
+                && success != failure)
+                return false;
+
+            // Memory scope without memory order is disallowed.
+            if (success == MEMORY_ORDER_EMPTY && scope != MEMORY_SCOPE_EMPTY)
+                return false;
+
+            // Failure must not be release or acq_rel.
+            if (failure == MEMORY_ORDER_RELEASE
+                || failure == MEMORY_ORDER_ACQ_REL)
+                return false;
+
+            // Failure must not be stronger than success.
+            return failure <= success;
+        };
+
         // repeat test for each reasonable memory order/scope combination
         std::vector<TExplicitMemoryOrderType> memoryOrder;
         std::vector<TExplicitMemoryScopeType> memoryScope;
@@ -719,16 +741,10 @@ class CBasicTestMemOrder2Scope
             {
                 for (unsigned si = 0; si < memoryScope.size(); si++)
                 {
-                    if ((memoryOrder[oi] == MEMORY_ORDER_EMPTY
-                         || memoryOrder[o2i] == MEMORY_ORDER_EMPTY)
-                        && memoryOrder[oi] != memoryOrder[o2i])
-                        continue; // both memory order arguments must be set (or
-                                  // none)
-                    if ((memoryOrder[oi] == MEMORY_ORDER_EMPTY
-                         || memoryOrder[o2i] == MEMORY_ORDER_EMPTY)
-                        && memoryScope[si] != MEMORY_SCOPE_EMPTY)
-                        continue; // memory scope without memory order is not
-                                  // allowed
+                    if (!checkValidity(memoryOrder[oi], memoryOrder[o2i],
+                                       memoryScope[si]))
+                        continue;
+
                     MemoryOrder(memoryOrder[oi]);
                     MemoryOrder2(memoryOrder[o2i]);
                     MemoryScope(memoryScope[si]);