diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 67faecf8..828d417f 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm32 new file mode 100644 index 00000000..66ed0350 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.5/non_uniform_broadcast_dynamic_index.spvasm64 new file mode 100644 index 00000000..f97d50d9 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm32 new file mode 100644 index 00000000..3707c432 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm64 new file mode 100644 index 00000000..76f38c41 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/test_spirv_15.cpp b/test_conformance/spirv_new/test_spirv_15.cpp new file mode 100644 index 00000000..dfab7e9a --- /dev/null +++ b/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 +#include +#include + +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(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 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(check)) + { + log_error("Result mismatch at index %zu! Got %d, Wanted %zu\n", + index, results[index], check); + return TEST_FAIL; + } + } + } + + return TEST_PASS; +}