diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 new file mode 100644 index 000000000..3076a0671 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 new file mode 100644 index 000000000..2e6247a82 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/copylogical_struct.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 3df78b455..091bd15b3 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/test_conformance/spirv_new/test_spirv_14.cpp @@ -511,3 +511,53 @@ TEST_SPIRV_FUNC(spirv14_select_composite) 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; +}