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 OpCopyLogical (#2136)

Browse Source 
This PR adds targeted testing for the SPIR-V 1.4 instruction
OpCopyLogical.
This commit is contained in:
Ben Ashbaugh
2024-11-25 07:35:50 -08:00
committed by GitHub
parent 2358b46553
commit e9a248f555
3 changed files with 96 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 Normal file
View File

@@ -0,0 +1,24 @@
; SPIR-V
; Version: 1.4
OpCapability Addresses
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %kernel "copylogical_test"
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%void = OpTypeVoid
%struct_a = OpTypeStruct %uint %float
%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a
%struct_b = OpTypeStruct %uint %float
%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b
%kernel_sig = OpTypeFunction %void %ptr_struct_b
%uint_1024 = OpConstant %uint 1024
%float_pi = OpConstant %float 3.1415
%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi
%kernel = OpFunction %void None %kernel_sig
%dst = OpFunctionParameter %ptr_struct_b
%entry = OpLabel
%struct_b_dst = OpCopyLogical %struct_b %struct_a_src
OpStore %dst %struct_b_dst
OpReturn
OpFunctionEnd

24
test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 Normal file
View File

@@ -0,0 +1,24 @@
; SPIR-V
; Version: 1.4
OpCapability Addresses
OpCapability Kernel
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %kernel "copylogical_test"
%uint = OpTypeInt 32 0
%float = OpTypeFloat 32
%void = OpTypeVoid
%struct_a = OpTypeStruct %uint %float
%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a
%struct_b = OpTypeStruct %uint %float
%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b
%kernel_sig = OpTypeFunction %void %ptr_struct_b
%uint_1024 = OpConstant %uint 1024
%float_pi = OpConstant %float 3.1415
%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi
%kernel = OpFunction %void None %kernel_sig
%dst = OpFunctionParameter %ptr_struct_b
%entry = OpLabel
%struct_b_dst = OpCopyLogical %struct_b %struct_a_src
OpStore %dst %struct_b_dst
OpReturn
OpFunctionEnd

48
test_conformance/spirv_new/test_spirv_14.cpp
View File

@@ -511,3 +511,51 @@ TEST_SPIRV_FUNC(spirv14_select_composite)
return TEST_PASS; return TEST_PASS;
} }
TEST_SPIRV_FUNC(spirv14_copylogical)
{
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;
}
cl_int error = CL_SUCCESS;
clProgramWrapper prog;
error = get_program_with_il(prog, deviceID, context,
"spv1.4/copylogical_struct");
SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
clKernelWrapper kernel = clCreateKernel(prog, "copylogical_test", &error);
SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
struct TestStruct
{
cl_int i;
cl_float f;
};
TestStruct results{ 0, 0.0f };
clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(results), NULL, &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(results),
&results, 0, NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
if (results.i != 1024 || results.f != 3.1415f)
{
log_error("Results mismatch! Got: { %d, %f }\n", results.i, results.f);
return TEST_FAIL;
}
return TEST_PASS;
}