From a256e4ad05afee25d0783a52fae86339d3c2efcd Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 6 Feb 2024 09:56:49 -0800 Subject: [PATCH] add test for cl_khr_spirv_linkonce_odr (#1226) * initial version of the test with placeholders for linkonce_odr linkage * add OpExtension SPV_KHR_linkonce_odr extension * add check for extension * switch to actual LinkOnceODR linkage * fix formatting * add a test case to ensure a function with linkonce_odr is exported * add back the extension check * fix formatting * undo compiler optimization and actually add the call to function a --- .../linkage_linkonce_odr_main.spvasm32 | 47 ++++++++++ .../linkage_linkonce_odr_main.spvasm64 | 51 ++++++++++ .../linkage_linkonce_odr_noa_main.spvasm32 | 44 +++++++++ .../linkage_linkonce_odr_noa_main.spvasm64 | 48 ++++++++++ .../linkage_linkonce_odr_obj.spvasm32 | 28 ++++++ .../linkage_linkonce_odr_obj.spvasm64 | 28 ++++++ test_conformance/spirv_new/test_linkage.cpp | 93 +++++++++++++++++++ 7 files changed, 339 insertions(+) create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm64 diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm32 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm32 new file mode 100644 index 00000000..dbdbe32e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm32 @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 27 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %17 "test_linkonce_odr" %__spirv_BuiltInGlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %18 FuncParamAttr NoCapture + OpDecorate %a LinkageAttributes "a" LinkOnceODR + OpDecorate %b LinkageAttributes "b" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %6 = OpTypeFunction %uint %uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %16 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input + %b = OpFunction %uint None %6 + %8 = OpFunctionParameter %uint + OpFunctionEnd + %a = OpFunction %uint Pure %6 + %10 = OpFunctionParameter %uint + %11 = OpLabel + %13 = OpIAdd %uint %10 %uint_5 + OpReturnValue %13 + OpFunctionEnd + %17 = OpFunction %void None %16 + %18 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %19 = OpLabel + %20 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %21 = OpCompositeExtract %uint %20 0 + %22 = OpFunctionCall %uint %a %21 + %23 = OpFunctionCall %uint %b %21 + %24 = OpIAdd %uint %22 %23 + %25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %18 %21 + OpStore %25 %24 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm64 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm64 new file mode 100644 index 00000000..243ab6b7 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm64 @@ -0,0 +1,51 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 30 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %18 "test_linkonce_odr" %__spirv_BuiltInGlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %19 FuncParamAttr NoCapture + OpDecorate %a LinkageAttributes "a" LinkOnceODR + OpDecorate %b LinkageAttributes "b" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %7 = OpTypeFunction %uint %uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %17 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input + %b = OpFunction %uint None %7 + %9 = OpFunctionParameter %uint + OpFunctionEnd + %a = OpFunction %uint Pure %7 + %11 = OpFunctionParameter %uint + %12 = OpLabel + %14 = OpIAdd %uint %11 %uint_5 + OpReturnValue %14 + OpFunctionEnd + %18 = OpFunction %void None %17 + %19 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %20 = OpLabel + %21 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32 + %22 = OpCompositeExtract %ulong %21 0 + %23 = OpUConvert %uint %22 + %24 = OpFunctionCall %uint %a %23 + %25 = OpFunctionCall %uint %b %23 + %26 = OpIAdd %uint %24 %25 + %27 = OpSConvert %ulong %23 + %28 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %19 %27 + OpStore %28 %26 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm32 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm32 new file mode 100644 index 00000000..e0b01b66 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm32 @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 27 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %17 "test_linkonce_odr" %__spirv_BuiltInGlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %18 FuncParamAttr NoCapture + OpDecorate %a LinkageAttributes "a" Import + OpDecorate %b LinkageAttributes "b" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %6 = OpTypeFunction %uint %uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %16 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input + %b = OpFunction %uint None %6 + %8 = OpFunctionParameter %uint + OpFunctionEnd + %a = OpFunction %uint None %6 + %10 = OpFunctionParameter %uint + OpFunctionEnd + %17 = OpFunction %void None %16 + %18 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %19 = OpLabel + %20 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %21 = OpCompositeExtract %uint %20 0 + %22 = OpFunctionCall %uint %a %21 + %23 = OpFunctionCall %uint %b %21 + %24 = OpIAdd %uint %22 %23 + %25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %18 %21 + OpStore %25 %24 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm64 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm64 new file mode 100644 index 00000000..bb022027 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_noa_main.spvasm64 @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 30 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %18 "test_linkonce_odr" %__spirv_BuiltInGlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %19 FuncParamAttr NoCapture + OpDecorate %a LinkageAttributes "a" Import + OpDecorate %b LinkageAttributes "b" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %7 = OpTypeFunction %uint %uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %17 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input + %b = OpFunction %uint None %7 + %9 = OpFunctionParameter %uint + OpFunctionEnd + %a = OpFunction %uint None %7 + %11 = OpFunctionParameter %uint + OpFunctionEnd + %18 = OpFunction %void None %17 + %19 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %20 = OpLabel + %21 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32 + %22 = OpCompositeExtract %ulong %21 0 + %23 = OpUConvert %uint %22 + %24 = OpFunctionCall %uint %a %23 + %25 = OpFunctionCall %uint %b %23 + %26 = OpIAdd %uint %24 %25 + %27 = OpSConvert %ulong %23 + %28 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %19 %27 + OpStore %28 %26 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm32 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm32 new file mode 100644 index 00000000..59c3f9d9 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm32 @@ -0,0 +1,28 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 14 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical32 OpenCL + OpDecorate %a LinkageAttributes "a" LinkOnceODR + OpDecorate %b LinkageAttributes "b" Export + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %uint_0 = OpConstant %uint 0 + %3 = OpTypeFunction %uint %uint + %a = OpFunction %uint Pure %3 + %5 = OpFunctionParameter %uint + %6 = OpLabel + %8 = OpIAdd %uint %5 %uint_5 + OpReturnValue %8 + OpFunctionEnd + %b = OpFunction %uint Pure %3 + %10 = OpFunctionParameter %uint + %11 = OpLabel + %13 = OpISub %uint %uint_0 %10 + OpReturnValue %13 + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm64 b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm64 new file mode 100644 index 00000000..5df6fdfe --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm64 @@ -0,0 +1,28 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 14 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpExtension "SPV_KHR_linkonce_odr" + OpMemoryModel Physical64 OpenCL + OpDecorate %a LinkageAttributes "a" LinkOnceODR + OpDecorate %b LinkageAttributes "b" Export + %uint = OpTypeInt 32 0 + %uint_5 = OpConstant %uint 5 + %uint_0 = OpConstant %uint 0 + %3 = OpTypeFunction %uint %uint + %a = OpFunction %uint Pure %3 + %5 = OpFunctionParameter %uint + %6 = OpLabel + %8 = OpIAdd %uint %5 %uint_5 + OpReturnValue %8 + OpFunctionEnd + %b = OpFunction %uint Pure %3 + %10 = OpFunctionParameter %uint + %11 = OpLabel + %13 = OpISub %uint %uint_0 %10 + OpReturnValue %13 + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_linkage.cpp b/test_conformance/spirv_new/test_linkage.cpp index c4635ee6..ea17040a 100644 --- a/test_conformance/spirv_new/test_linkage.cpp +++ b/test_conformance/spirv_new/test_linkage.cpp @@ -144,3 +144,96 @@ TEST_SPIRV_FUNC(linkage_import_function_link) return 0; } + +static int test_linkonce_odr_helper(cl_device_id deviceID, cl_context context, + cl_command_queue queue, + const char *main_module_filename) +{ + cl_int err = 0; + + clProgramWrapper prog_obj; + err = test_linkage_compile(deviceID, context, queue, + "linkage_linkonce_odr_obj", prog_obj); + SPIRV_CHECK_ERROR(err, "Failed to compile export program"); + + clProgramWrapper prog_main; + err = test_linkage_compile(deviceID, context, queue, main_module_filename, + prog_main); + SPIRV_CHECK_ERROR(err, "Failed to compile import program"); + + cl_program progs[] = { prog_obj, prog_main }; + + clProgramWrapper prog = + clLinkProgram(context, 1, &deviceID, NULL, 2, progs, NULL, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to link programs"); + + clKernelWrapper kernel = clCreateKernel(prog, "test_linkonce_odr", &err); + SPIRV_CHECK_ERROR(err, "Failed to create spv kernel"); + + const int num = 256; + std::vector h_in(num); + RandomSeed seed(gRandomSeed); + for (int i = 0; i < num; i++) + { + h_in[i] = genrand(seed) % 2048; + } + + size_t bytes = sizeof(cl_int) * num; + clMemWrapper in = + clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); + SPIRV_CHECK_ERROR(err, "Failed to create in buffer"); + + err = clEnqueueWriteBuffer(queue, in, CL_TRUE, 0, bytes, &h_in[0], 0, NULL, + NULL); + SPIRV_CHECK_ERROR(err, "Failed to copy to in buffer"); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &in); + SPIRV_CHECK_ERROR(err, "Failed to set arg 1"); + + size_t global = num; + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, + NULL); + SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel"); + + std::vector h_out(num); + err = clEnqueueReadBuffer(queue, in, CL_TRUE, 0, bytes, &h_out[0], 0, NULL, + NULL); + SPIRV_CHECK_ERROR(err, "Failed to read to output"); + + for (int i = 0; i < num; i++) + { + if (h_out[i] != 5) + { + log_error("Incorrect values at location %d\n", i); + return TEST_FAIL; + } + } + + return TEST_PASS; +} + +TEST_SPIRV_FUNC(linkage_linkonce_odr) +{ + if (!is_extension_available(deviceID, "cl_khr_spirv_linkonce_odr")) + { + log_info("Extension cl_khr_spirv_linkonce_odr not supported; skipping " + "tests.\n"); + return TEST_SKIPPED_ITSELF; + } + + int result = TEST_PASS; + + // For this test, use the default main module, which has an "a" function + // with the linkonce_odr linkage type. This ensures that having two "a" + // functions with linkonce_odr works properly. + result |= test_linkonce_odr_helper(deviceID, context, queue, + "linkage_linkonce_odr_main"); + + // For this test, use a main module without the "a" function. This ensures + // that the "a" function is properly exported with the linkonce_odr linkage + // type. + result |= test_linkonce_odr_helper(deviceID, context, queue, + "linkage_linkonce_odr_noa_main"); + + return result; +}