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
2026-03-19 06:09:01 +00:00 · 2024-02-06 09:56:49 -08:00
parent d394b0f274
commit a256e4ad05
7 changed files with 339 additions and 0 deletions
--- a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm32
+++ 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
--- a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_main.spvasm64
+++ 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
--- 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
@@ -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
--- 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
@@ -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
--- a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm32
+++ 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
--- a/test_conformance/spirv_new/spirv_asm/linkage_linkonce_odr_obj.spvasm64
+++ 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
--- 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<cl_int> h_in(num);
+    RandomSeed seed(gRandomSeed);
+    for (int i = 0; i < num; i++)
+    {
+        h_in[i] = genrand<cl_int>(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<cl_int> 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;
+}