add testing for SPIR-V 1.5 (#2208)

fixes #2140 

Adds testing for SPIR-V 1.5 features:

* Adds a test for bitcasts between pointers and vectors of integers.
Note, SPIR-V 1.5 only supports bitcasts to vectors of two 32-bit
integers. Therefore, the SPIR-V 1.5 behavior will only be exercised on
devices with 64-bit pointers. The test will run on devices with 32-bit
pointers, but will instead bitcast to scalars.
* Adds a test for OpGroupNonUniformBroadcast with a dynamic index. Note,
this is not an exhaustive test, and only unsigned integer types are
tested, to avoid duplicating testing for cl_khr_subgroup_ballot.

test_conformance/spirv_new/CMakeLists.txt

@@ -28,6 +28,7 @@ set(${MODULE_NAME}_SOURCES
   test_op_vector_insert.cpp
   test_op_vector_times_scalar.cpp
   test_spirv_14.cpp
+  test_spirv_15.cpp
 )
 
 set(TEST_HARNESS_SOURCES

test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm32

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.5
+; Reference:
+;   kernel void non_uniform_broadcast_dynamic_index_test(global uint* dst_base) {
+;       uint id = get_global_id(0);
+;       uint index = get_group_id(0);
+;       uint value = sub_group_non_uniform_broadcast(id, index);
+;       dst_base[id] = value;
+;   }
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability GroupNonUniformBallot
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "non_uniform_broadcast_dynamic_index_test" %pglobalid %pgroupid
+               OpDecorate %pglobalid BuiltIn GlobalInvocationId
+               OpDecorate %pgroupid BuiltIn WorkgroupId
+       %uint = OpTypeInt 32 0
+   %sg_scope = OpConstant %uint 3
+      %uint3 = OpTypeVector %uint 3
+       %void = OpTypeVoid
+%iptr_uint3 = OpTypePointer Input %uint3
+  %gptr_uint = OpTypePointer CrossWorkgroup %uint
+ %kernel_sig = OpTypeFunction %void %gptr_uint
+  %pglobalid = OpVariable %iptr_uint3 Input
+   %pgroupid = OpVariable %iptr_uint3 Input
+     %kernel = OpFunction %void None %kernel_sig
+   %dst_base = OpFunctionParameter %gptr_uint
+      %entry = OpLabel
+   %globalid = OpLoad %uint3 %pglobalid Aligned 32
+         %id = OpCompositeExtract %uint %globalid 0
+    %groupid = OpLoad %uint3 %pgroupid Aligned 32
+      %index = OpCompositeExtract %uint %groupid 0
+      %value = OpGroupNonUniformBroadcast %uint %sg_scope %id %index
+        %dst = OpInBoundsPtrAccessChain %gptr_uint %dst_base %id
+               OpStore %dst %value
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm64

@@ -0,0 +1,41 @@
+; SPIR-V
+; Version: 1.5
+; Reference:
+;   kernel void non_uniform_broadcast_dynamic_index_test(global uint* dst_base) {
+;       uint id = get_global_id(0);
+;       uint index = get_group_id(0);
+;       uint value = sub_group_non_uniform_broadcast(id, index);
+;       dst_base[id] = value;
+;   }
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpCapability GroupNonUniformBallot
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "non_uniform_broadcast_dynamic_index_test" %pglobalid %pgroupid
+               OpDecorate %pglobalid BuiltIn GlobalInvocationId
+               OpDecorate %pgroupid BuiltIn WorkgroupId
+       %uint = OpTypeInt 32 0
+   %sg_scope = OpConstant %uint 3
+      %ulong = OpTypeInt 64 0
+     %ulong3 = OpTypeVector %ulong 3
+       %void = OpTypeVoid
+%iptr_ulong3 = OpTypePointer Input %ulong3
+  %gptr_uint = OpTypePointer CrossWorkgroup %uint
+ %kernel_sig = OpTypeFunction %void %gptr_uint
+  %pglobalid = OpVariable %iptr_ulong3 Input
+   %pgroupid = OpVariable %iptr_ulong3 Input
+     %kernel = OpFunction %void None %kernel_sig
+   %dst_base = OpFunctionParameter %gptr_uint
+      %entry = OpLabel
+   %globalid = OpLoad %ulong3 %pglobalid Aligned 32
+  %globalid0 = OpCompositeExtract %ulong %globalid 0
+         %id = OpUConvert %uint %globalid0
+    %groupid = OpLoad %ulong3 %pgroupid Aligned 32
+   %groupid0 = OpCompositeExtract %ulong %groupid 0
+      %index = OpUConvert %uint %groupid0
+      %value = OpGroupNonUniformBroadcast %uint %sg_scope %id %index
+        %dst = OpInBoundsPtrAccessChain %gptr_uint %dst_base %globalid0
+               OpStore %dst %value
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm32

@@ -0,0 +1,22 @@
+; SPIR-V
+; Version: 1.5
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "ptr_bitcast_test"
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+   %pptr_int = OpTypePointer Function %uint
+  %gptr_uint = OpTypePointer CrossWorkgroup %uint
+ %kernel_sig = OpTypeFunction %void %gptr_uint %gptr_uint
+    %uint_42 = OpConstant %uint 42
+     %kernel = OpFunction %void None %kernel_sig
+  %dst_uint0 = OpFunctionParameter %gptr_uint
+  %dst_uint1 = OpFunctionParameter %gptr_uint
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+   %uint_ptr = OpBitcast %uint %pvalue
+               OpStore %dst_uint0 %uint_ptr
+               OpStore %dst_uint1 %uint_ptr
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm64

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.5
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "ptr_bitcast_test"
+       %uint = OpTypeInt 32 0
+      %ulong = OpTypeInt 64 0
+      %uint2 = OpTypeVector %uint 2
+       %void = OpTypeVoid
+   %pptr_int = OpTypePointer Function %uint
+ %gptr_ulong = OpTypePointer CrossWorkgroup %ulong
+ %gptr_uint2 = OpTypePointer CrossWorkgroup %uint2
+ %kernel_sig = OpTypeFunction %void %gptr_ulong %gptr_uint2
+    %uint_42 = OpConstant %uint 42
+     %kernel = OpFunction %void None %kernel_sig
+  %dst_ulong = OpFunctionParameter %gptr_ulong
+  %dst_uint2 = OpFunctionParameter %gptr_uint2
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+  %ulong_ptr = OpBitcast %ulong %pvalue
+               OpStore %dst_ulong %ulong_ptr
+  %uint2_ptr = OpBitcast %uint2 %pvalue
+               OpStore %dst_uint2 %uint2_ptr
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/test_spirv_15.cpp

@@ -0,0 +1,162 @@
+//
+// 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.
+// 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 "spirvInfo.hpp"
+#include "types.hpp"
+
+#include <algorithm>
+#include <cinttypes>
+#include <vector>
+
+REGISTER_TEST(spirv15_ptr_bitcast)
+{
+    if (!is_spirv_version_supported(device, "SPIR-V_1.5"))
+    {
+        log_info("SPIR-V 1.5 not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    cl_uint address_bits;
+    error = clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint),
+                            &address_bits, NULL);
+    SPIRV_CHECK_ERROR(error, "Failed to get address bits");
+
+    clProgramWrapper prog;
+    error = get_program_with_il(prog, device, context, "spv1.5/ptr_bitcast");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "ptr_bitcast_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    cl_ulong result_ulong =
+        address_bits == 32 ? 0xAAAAAAAAUL : 0xAAAAAAAAAAAAAAAAUL;
+    cl_ulong result_uint2 =
+        address_bits == 32 ? 0x55555555UL : 0x5555555555555555UL;
+
+    clMemWrapper dst_ulong =
+        clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                       sizeof(result_ulong), &result_ulong, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst_ulong buffer");
+
+    clMemWrapper dst_uint2 =
+        clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                       sizeof(result_uint2), &result_uint2, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst_uint2 buffer");
+
+    error |= clSetKernelArg(kernel, 0, sizeof(dst_ulong), &dst_ulong);
+    error |= clSetKernelArg(kernel, 1, sizeof(dst_uint2), &dst_uint2);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    size_t global = 1;
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0,
+                                   NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    error =
+        clEnqueueReadBuffer(queue, dst_ulong, CL_TRUE, 0, sizeof(result_ulong),
+                            &result_ulong, 0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read dst_ulong buffer");
+
+    error =
+        clEnqueueReadBuffer(queue, dst_uint2, CL_TRUE, 0, sizeof(result_uint2),
+                            &result_uint2, 0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read dst_uint2 buffer");
+
+    if (result_ulong != result_uint2)
+    {
+        log_error("Results mismatch!  ulong = 0x%016" PRIx64
+                  " vs. uint2 = 0x%016" PRIx64 "\n",
+                  result_ulong, result_uint2);
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}
+
+REGISTER_TEST(spirv15_non_uniform_broadcast)
+{
+    if (!is_spirv_version_supported(device, "SPIR-V_1.5"))
+    {
+        log_info("SPIR-V 1.5 not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    if (!is_extension_available(device, "cl_khr_subgroup_ballot"))
+    {
+        log_info("cl_khr_subgroup_ballot is not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    clProgramWrapper prog;
+    error = get_program_with_il(prog, device, context,
+                                "spv1.5/non_uniform_broadcast_dynamic_index");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(
+        prog, "non_uniform_broadcast_dynamic_index_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    // Get the local work-group size for one sub-group per work-group.
+    size_t lws = 0;
+    size_t one = 1;
+    error = clGetKernelSubGroupInfo(
+        kernel, device, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
+        sizeof(size_t), &one, sizeof(size_t), &lws, NULL);
+    SPIRV_CHECK_ERROR(error, "Failed to get local work size for one sub-group");
+
+    // Use four work-groups, unless the local-group size is less than four.
+    size_t wgcount = std::min<size_t>(lws, 4);
+    size_t gws = wgcount * lws;
+    clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                      sizeof(cl_int) * gws, NULL, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
+
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &gws, &lws, 0, NULL,
+                                   NULL);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    std::vector<cl_int> results(gws);
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(cl_int) * gws,
+                                results.data(), 0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    // Remember: the test kernel did:
+    //  sub_group_non_uniform_broadcast(get_global_id(0), get_group_id(0))
+    for (size_t g = 0; g < wgcount; g++)
+    {
+        for (size_t l = 0; l < lws; l++)
+        {
+            size_t index = g * lws + l;
+            size_t check = g * lws + g;
+            if (results[index] != static_cast<cl_int>(check))
+            {
+                log_error("Result mismatch at index %zu!  Got %d, Wanted %zu\n",
+                          index, results[index], check);
+                return TEST_FAIL;
+            }
+        }
+    }
+
+    return TEST_PASS;
+}