mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
add SPIR-V 1.6 testing for UniformDecoration decorations (#2254)
Adds a basic test for the SPIR-V 1.6 UniformDecoration decorations. Specifically: * Tests both the Uniform and UniformId decorations. * Tests the decorations on constants, function parameters, and variables.
This commit is contained in:
@@ -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
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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<cl_uint> 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);
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user