From c4119f07c1aeb29b4ae086a6d4e674169b846e0c Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Thu, 31 Oct 2024 07:12:30 -0700 Subject: [PATCH] add SPIR-V 1.4 testing for UserSemantic (#2053) This PR adds targeted testing for more SPIR-V 1.4 features. Specifically, this PR adds testing for OpDecorateString, OpMemberDecorateString, and the UserSemantic decoration. --- .../usersemantic_decoratestring.spvasm32 | 35 ++++++++++ .../usersemantic_decoratestring.spvasm64 | 39 +++++++++++ ...usersemantic_memberdecoratestring.spvasm32 | 39 +++++++++++ ...usersemantic_memberdecoratestring.spvasm64 | 43 ++++++++++++ test_conformance/spirv_new/test_spirv_14.cpp | 66 +++++++++++++++++++ 5 files changed, 222 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm64 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm32 new file mode 100644 index 0000000000..9a85f39ce4 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm32 @@ -0,0 +1,35 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 24 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %usersemantic_test "usersemantic_test" %global_id + OpDecorate %global_id LinkageAttributes "global_id" Import + OpDecorate %global_id Constant + OpDecorate %global_id BuiltIn GlobalInvocationId + ; Basic decoration: + OpDecorateString %global_id UserSemantic "FOO" + ; Duplicate decorations are allowed as long as the string is different. + OpDecorateString %global_id UserSemantic "BAR" + ; Try one more string with punctuation. + OpDecorateString %global_id UserSemantic "FOO? BAR. BAZ!" + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%global_id = OpVariable %_ptr_Input_v3uint Input +%usersemantic_test = OpFunction %void None %9 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %index = OpLoad %v3uint %global_id Aligned 32 + %call = OpCompositeExtract %uint %index 0 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %call + OpStore %arrayidx %call Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm64 new file mode 100644 index 0000000000..b97b46af22 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm64 @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 24 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %usersemantic_test "usersemantic_test" %global_id + OpDecorate %global_id LinkageAttributes "global_id" Import + OpDecorate %global_id Constant + OpDecorate %global_id BuiltIn GlobalInvocationId + ; Basic decoration: + OpDecorateString %global_id UserSemantic "FOO" + ; Duplicate decorations are allowed as long as the string is different. + OpDecorateString %global_id UserSemantic "BAR" + ; Try one more string with punctuation. + OpDecorateString %global_id UserSemantic "FOO? BAR. BAZ!" + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%global_id = OpVariable %_ptr_Input_v3ulong Input +%usersemantic_test = OpFunction %void None %9 + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %index = OpLoad %v3ulong %global_id Aligned 32 + %call = OpCompositeExtract %ulong %index 0 + %conv = OpUConvert %uint %call + %idxprom = OpSConvert %ulong %conv + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %idxprom + OpStore %arrayidx %conv Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm32 new file mode 100644 index 0000000000..8b9a7ecfbc --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm32 @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 60 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %kernel "usersemantic_test" %global_id + OpDecorate %global_id LinkageAttributes "global_id" Import + OpDecorate %global_id Constant + OpDecorate %global_id BuiltIn GlobalInvocationId + OpDecorate %s Alignment 4 + OpMemberDecorateString %struct 0 UserSemantic "foo" + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%struct = OpTypeStruct %uint +%_ptr_Function_struct = OpTypePointer Function %struct +%_ptr_Function_uint = OpTypePointer Function %uint + %global_id = OpVariable %_ptr_Input_v3uint Input + %kernel = OpFunction %void None %kernel_sig + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %s = OpVariable %_ptr_Function_struct Function + %gid = OpLoad %v3uint %global_id Aligned 16 + %gid0 = OpCompositeExtract %uint %gid 0 + %x = OpInBoundsPtrAccessChain %_ptr_Function_uint %s %uint_0 %uint_0 + OpStore %x %gid0 Aligned 4 + %index = OpLoad %uint %x Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %index + OpStore %arrayidx %index Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm64 new file mode 100644 index 0000000000..7138c8578f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm64 @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 60 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %kernel "usersemantic_test" %global_id + OpDecorate %global_id LinkageAttributes "global_id" Import + OpDecorate %global_id Constant + OpDecorate %global_id BuiltIn GlobalInvocationId + OpDecorate %s Alignment 4 + OpMemberDecorateString %struct 0 UserSemantic "foo" + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %kernel_sig = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%struct = OpTypeStruct %uint +%_ptr_Function_struct = OpTypePointer Function %struct +%_ptr_Function_uint = OpTypePointer Function %uint + %global_id = OpVariable %_ptr_Input_v3ulong Input + %kernel = OpFunction %void None %kernel_sig + %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %s = OpVariable %_ptr_Function_struct Function + %gid = OpLoad %v3ulong %global_id Aligned 32 + %gid0 = OpCompositeExtract %ulong %gid 0 + %conv = OpUConvert %uint %gid0 + %x = OpInBoundsPtrAccessChain %_ptr_Function_uint %s %uint_0 %uint_0 + OpStore %x %conv Aligned 4 + %index = OpLoad %uint %x Aligned 4 + %idxprom = OpSConvert %ulong %index + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %idxprom + OpStore %arrayidx %index Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 6c6b890f4f..60d74e57a2 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -286,3 +286,69 @@ TEST_SPIRV_FUNC(spirv14_ptrops) return TEST_PASS; } + +static int test_usersemantic_decoration(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + bool test_memberdecoratestring) +{ + cl_int error = CL_SUCCESS; + + const char* filename = test_memberdecoratestring + ? "spv1.4/usersemantic_memberdecoratestring" + : "spv1.4/usersemantic_decoratestring"; + + clProgramWrapper prog; + error = get_program_with_il(prog, deviceID, context, filename); + SPIRV_CHECK_ERROR(error, "Failed to compile spv program"); + + clKernelWrapper kernel = clCreateKernel(prog, "usersemantic_test", &error); + SPIRV_CHECK_ERROR(error, "Failed to create spv kernel"); + + int h_dst = -1; + clMemWrapper dst = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + sizeof(h_dst), &h_dst, &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"); + + 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, CL_TRUE, 0, sizeof(h_dst), &h_dst, + 0, NULL, NULL); + SPIRV_CHECK_ERROR(error, "Unable to read destination buffer"); + + if (h_dst != 0) + { + log_error("Mismatch! Got: %i, Wanted: %i\n", h_dst, 0); + return TEST_FAIL; + } + + return TEST_PASS; +} + +TEST_SPIRV_FUNC(spirv14_usersemantic_decoratestring) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_usersemantic_decoration(deviceID, context, queue, false); +} + +TEST_SPIRV_FUNC(spirv14_usersemantic_memberdecoratestring) +{ + if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) + { + log_info("SPIR-V 1.4 not supported; skipping tests.\n"); + return TEST_SKIPPED_ITSELF; + } + + return test_usersemantic_decoration(deviceID, context, queue, true); +}