ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

add SPIR-V 1.4 testing for UserSemantic (#2053)

Browse Source 
This PR adds targeted testing for more SPIR-V 1.4 features.
Specifically, this PR adds testing for OpDecorateString,
OpMemberDecorateString, and the UserSemantic decoration.
This commit is contained in:
Ben Ashbaugh
2024-10-31 07:12:30 -07:00
committed by GitHub
parent d8228f0d72
commit c4119f07c1
5 changed files with 222 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

35
test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm32 Normal file
View File

@@ -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

39
test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_decoratestring.spvasm64 Normal file
View File

@@ -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

39
test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm32 Normal file
View File

@@ -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

43
test_conformance/spirv_new/spirv_asm/spv1.4/usersemantic_memberdecoratestring.spvasm64 Normal file
View File

@@ -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

66
test_conformance/spirv_new/test_spirv_14.cpp
View File

@@ -286,3 +286,69 @@ TEST_SPIRV_FUNC(spirv14_ptrops)
return TEST_PASS; 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);
}