[SYCL] Allow an inlined kernel to be called from multiple TUs.

clang/lib/AST/ASTContext.cpp

@@ -9729,6 +9729,10 @@ static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
     // visible externally so they can be launched from host.
     if (L == GVA_DiscardableODR || L == GVA_Internal)
       return GVA_StrongODR;
+  } else if (Context.getLangOpts().SYCLIsDevice &&
+             D->hasAttr<OpenCLKernelAttr>()) {
+    if (L == GVA_DiscardableODR)
+      return GVA_StrongODR;
   }
   return L;
 }

clang/lib/Sema/SemaSYCL.cpp

@@ -948,6 +948,11 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc) {
   // Let's copy source location of a functor/lambda to emit nicer diagnostics
   OpenCLKernel->setLocation(LE->getLocation());
 
+  // If the source function is implicitly inline, the kernel should be marked
+  // such as well. This allows the kernel to be ODR'd if there are multiple uses
+  // in different translation units.
+  OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined());
+
   CompoundStmt *OpenCLKernelBody =
       CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel);
   OpenCLKernel->setBody(OpenCLKernelBody);

clang/test/CodeGenSYCL/image_accessor.cpp

@@ -7,27 +7,27 @@
 // RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
 //
 // CHECK-1DRO: %opencl.image1d_ro_t = type opaque
-// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-1DRO: define weak_odr spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-2DRO: %opencl.image2d_ro_t = type opaque
-// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-2DRO: define weak_odr spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DRO: %opencl.image3d_ro_t = type opaque
-// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-3DRO: define weak_odr spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-1DWO: %opencl.image1d_wo_t = type opaque
-// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-1DWO: define weak_odr spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-2DWO: %opencl.image2d_wo_t = type opaque
-// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-2DWO: define weak_odr spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DWO: %opencl.image3d_wo_t = type opaque
-// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-3DWO: define weak_odr spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // TODO: Add tests for the image_array opencl datatype support.

sycl/test/separate-compile/same-kernel.cpp

@@ -0,0 +1,73 @@
+//==----------- same-kernel.cpp --------------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// >> ---- compile src1
+// RUN: %clang -std=c++11 -fsycl -c %s -o %t-same-kernel-a.o
+//
+// >> ---- compile src2
+// RUN: %clang -DB_CPP=1 -std=c++11 -fsycl -c %s -o %t-same-kernel-b.o
+//
+// >> ---- link the full hetero app
+// RUN: %clang %t-same-kernel-a.o %t-same-kernel-b.o -o %t-same-kernel.exe -fsycl -lstdc++
+// RUN: %CPU_RUN_PLACEHOLDER %t-same-kernel.exe
+// RUN: %GPU_RUN_PLACEHOLDER %t-same-kernel.exe
+// RUN: %ACC_RUN_PLACEHOLDER %t-same-kernel.exe
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+class TestFnObj {
+public:
+  TestFnObj(buffer<int> &buf, handler &cgh) :
+    data(buf.get_access<access::mode::write>(cgh)) {}
+  accessor<int, 1, access::mode::write, access::target::global_buffer> data;
+  void operator()(id<1> item) {
+    data[item] = item[0];
+  }
+};
+
+void kernel2();
+
+#ifndef B_CPP
+void kernel2() {
+  static int data[256];
+  {
+    buffer<int> b(data, range<1>(256));
+    queue q;
+    q.submit([&](handler &cgh){
+      TestFnObj kernel(b, cgh);
+      cgh.parallel_for(range<1>(256), kernel);
+    });
+  }
+  for (int i = 0; i < 256; i++) {
+    assert(data[i] == i);
+  }
+}
+#else // B_CPP
+void kernel1() {
+  static int data[10];
+  {
+    buffer<int> b(data, range<1>(10));
+    queue q;
+    q.submit([&](cl::sycl::handler &cgh){
+      TestFnObj kernel(b, cgh);
+      cgh.parallel_for(range<1>(10), kernel);
+    });
+  }
+  for (int i = 0; i < 10; i++) {
+    assert(data[i] == i);
+  }
+}
+
+int main() {
+  kernel1();
+  kernel2();
+
+  return 0;
+}
+#endif // B_CPP