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.
KhronosGroup · Oct 31, 2024 · c4119f0 · c4119f0
1 parent d8228f0
commit c4119f0
Show file tree

Hide file tree

Showing 5 changed files with 222 additions and 0 deletions.
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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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);
+}