diff --git a/test_conformance/spirv_new/spirv_asm/CMakeLists.txt b/test_conformance/spirv_new/spirv_asm/CMakeLists.txt index 0e12a285..34bb30bf 100644 --- a/test_conformance/spirv_new/spirv_asm/CMakeLists.txt +++ b/test_conformance/spirv_new/spirv_asm/CMakeLists.txt @@ -411,6 +411,12 @@ set(spirv_sources spv1.5/basic.spvasm64 spv1.6/basic.spvasm32 spv1.6/basic.spvasm64 + spv1.6/image_operand_nontemporal.spvasm32 + spv1.6/image_operand_nontemporal.spvasm64 + spv1.6/uniformdecoration_uniformid.spvasm32 + spv1.6/uniformdecoration_uniformid.spvasm64 + spv1.6/uniformdecoration_uniform.spvasm32 + spv1.6/uniformdecoration_uniform.spvasm64 undef_char_simple.spvasm32 undef_char_simple.spvasm64 undef_double_simple.spvasm32 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm32 new file mode 100644 index 00000000..be2c0143 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm32 @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.6 + OpCapability Addresses + OpCapability Kernel + OpCapability UniformDecoration + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %kernel "test_uniformdecoration" %gid + OpDecorate %gid BuiltIn GlobalInvocationId + OpDecorate %gid Constant + ; Decoration on a constant + OpDecorate %uint_0 Uniform + ; Decoration on a function parameter + OpDecorate %value Uniform + OpDecorate %base Uniform + ; Decoration on a variable + OpDecorate %newvalue Uniform + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %ptr_gid = OpTypePointer Input %v3uint + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %gptr_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %gptr_uint %uint + %gid = OpVariable %ptr_gid Input + %kernel = OpFunction %void None %kernel_sig + %base = OpFunctionParameter %gptr_uint + %value = OpFunctionParameter %uint + %entry = OpLabel + %gidv3 = OpLoad %v3uint %gid Aligned 32 + %gid0 = OpCompositeExtract %uint %gidv3 0 + %ptr = OpInBoundsPtrAccessChain %gptr_uint %base %gid0 + %newvalue = OpIAdd %uint %value %uint_1 + OpStore %ptr %newvalue Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm64 new file mode 100644 index 00000000..6cefca8c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniform.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.6 + OpCapability Addresses + OpCapability Kernel + OpCapability Int64 + OpCapability UniformDecoration + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %kernel "test_uniformdecoration" %gid + OpDecorate %gid BuiltIn GlobalInvocationId + OpDecorate %gid Constant + ; Decoration on a constant + OpDecorate %uint_0 Uniform + ; Decoration on a function parameter + OpDecorate %value Uniform + OpDecorate %base Uniform + ; Decoration on a variable + OpDecorate %newvalue Uniform + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %v3ulong = OpTypeVector %ulong 3 + %ptr_gid = OpTypePointer Input %v3ulong + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %gptr_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %gptr_uint %uint + %gid = OpVariable %ptr_gid Input + %kernel = OpFunction %void None %kernel_sig + %base = OpFunctionParameter %gptr_uint + %value = OpFunctionParameter %uint + %entry = OpLabel + %gidv3 = OpLoad %v3ulong %gid Aligned 32 + %gid0 = OpCompositeExtract %ulong %gidv3 0 + %ptr = OpInBoundsPtrAccessChain %gptr_uint %base %gid0 + %newvalue = OpIAdd %uint %value %uint_1 + OpStore %ptr %newvalue Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm32 new file mode 100644 index 00000000..db4c23a4 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm32 @@ -0,0 +1,40 @@ +; SPIR-V +; Version: 1.6 + OpCapability Addresses + OpCapability Kernel + OpCapability UniformDecoration + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %kernel "test_uniformdecoration" %gid + OpDecorate %gid BuiltIn GlobalInvocationId + OpDecorate %gid Constant + ; Decoration on a constant + OpDecorateId %uint_0 UniformId %scope_CrossDevice + ; Decoration on a function parameter + OpDecorateId %value UniformId %scope_Device + OpDecorateId %base UniformId %scope_Workgroup + ; Decoration on a variable + OpDecorateId %newvalue UniformId %scope_Subgroup + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %ptr_gid = OpTypePointer Input %v3uint +%scope_CrossDevice = OpConstant %uint 0 +%scope_Device = OpConstant %uint 1 +%scope_Workgroup = OpConstant %uint 2 +%scope_Subgroup = OpConstant %uint 3 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %gptr_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %gptr_uint %uint + %gid = OpVariable %ptr_gid Input + %kernel = OpFunction %void None %kernel_sig + %base = OpFunctionParameter %gptr_uint + %value = OpFunctionParameter %uint + %entry = OpLabel + %gidv3 = OpLoad %v3uint %gid Aligned 32 + %gid0 = OpCompositeExtract %uint %gidv3 0 + %ptr = OpInBoundsPtrAccessChain %gptr_uint %base %gid0 + %newvalue = OpIAdd %uint %value %uint_1 + OpStore %ptr %newvalue Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm64 new file mode 100644 index 00000000..156ed536 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/uniformdecoration_uniformid.spvasm64 @@ -0,0 +1,42 @@ +; SPIR-V +; Version: 1.6 + OpCapability Addresses + OpCapability Kernel + OpCapability Int64 + OpCapability UniformDecoration + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %kernel "test_uniformdecoration" %gid + OpDecorate %gid BuiltIn GlobalInvocationId + OpDecorate %gid Constant + ; Decoration on a constant + OpDecorateId %uint_0 UniformId %scope_CrossDevice + ; Decoration on a function parameter + OpDecorateId %value UniformId %scope_Device + OpDecorateId %base UniformId %scope_Workgroup + ; Decoration on a variable + OpDecorateId %newvalue UniformId %scope_Subgroup + %uint = OpTypeInt 32 0 + %ulong = OpTypeInt 64 0 + %v3ulong = OpTypeVector %ulong 3 + %ptr_gid = OpTypePointer Input %v3ulong +%scope_CrossDevice = OpConstant %uint 0 +%scope_Device = OpConstant %uint 1 +%scope_Workgroup = OpConstant %uint 2 +%scope_Subgroup = OpConstant %uint 3 + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %gptr_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %gptr_uint %uint + %gid = OpVariable %ptr_gid Input + %kernel = OpFunction %void None %kernel_sig + %base = OpFunctionParameter %gptr_uint + %value = OpFunctionParameter %uint + %entry = OpLabel + %gidv3 = OpLoad %v3ulong %gid Aligned 32 + %gid0 = OpCompositeExtract %ulong %gidv3 0 + %ptr = OpInBoundsPtrAccessChain %gptr_uint %base %gid0 + %newvalue = OpIAdd %uint %value %uint_1 + OpStore %ptr %newvalue Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_spirv_16.cpp b/test_conformance/spirv_new/test_spirv_16.cpp index 169a2cb7..9c9b05a5 100644 --- a/test_conformance/spirv_new/test_spirv_16.cpp +++ b/test_conformance/spirv_new/test_spirv_16.cpp @@ -82,3 +82,79 @@ REGISTER_TEST(spirv16_image_operand_nontemporal) return TEST_PASS; } + +static int test_uniformdecoration_helper(cl_device_id device, + cl_context context, + cl_command_queue queue, + bool test_uniformid) +{ + constexpr size_t global_size = 16; + const cl_uint value = 42; + const cl_uint check = value + 1; + + const char* filename = test_uniformid ? "spv1.6/uniformdecoration_uniformid" + : "spv1.6/uniformdecoration_uniform"; + + cl_int error = CL_SUCCESS; + + clProgramWrapper prog; + error = get_program_with_il(prog, device, context, filename); + SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + + clKernelWrapper kernel = + clCreateKernel(prog, "test_uniformdecoration", &error); + SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + + std::vector h_dst(global_size); + clMemWrapper dst = + clCreateBuffer(context, CL_MEM_READ_WRITE, + h_dst.size() * sizeof(cl_uint), nullptr, &error); + SPIRV_CHECK_ERROR(error, "Failed to create dst buffer"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(value), &value); + SPIRV_CHECK_ERROR(error, "Failed to set kernel args"); + + error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_size, + nullptr, 0, nullptr, nullptr); + SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel"); + + error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, + h_dst.size() * sizeof(cl_uint), h_dst.data(), 0, + nullptr, nullptr); + SPIRV_CHECK_ERROR(error, "Unable to read dst buffer"); + + for (size_t i = 0; i < global_size; i++) + { + if (h_dst[i] != check) + { + log_error("Result mismatch at index %zu! Got %u, wanted %u.\n", i, + h_dst[i], check); + return TEST_FAIL; + } + } + + return TEST_PASS; +} + +REGISTER_TEST(spirv16_uniformdecoration_uniform) +{ + if (!is_spirv_version_supported(device, "SPIR-V_1.6")) + { + log_info("SPIR-V 1.6 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_uniformdecoration_helper(device, context, queue, false); +} + +REGISTER_TEST(spirv16_uniformdecoration_uniformid) +{ + if (!is_spirv_version_supported(device, "SPIR-V_1.6")) + { + log_info("SPIR-V 1.6 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_uniformdecoration_helper(device, context, queue, true); +}