[PATCH] D90311: [CUDA][HIP] Fix linkage for -fgpu-rdc

[Yaxun Liu via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGabd8cd9199d1: [CUDA][HIP] Fix linkage for -fgpu-rdc (authored by yaxunl).
Herald added a project: clang.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D90311/new/

https://reviews.llvm.org/D90311

Files:
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/CodeGenCUDA/device-fun-linkage.cu


Index: clang/test/CodeGenCUDA/device-fun-linkage.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/device-fun-linkage.cu
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
+// RUN:   -emit-llvm -o - %s \
+// RUN:   | FileCheck -check-prefix=NORDC %s
+// RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
+// RUN:   -fgpu-rdc -emit-llvm -o - %s \
+// RUN:   | FileCheck -check-prefix=RDC %s
+
+#include "Inputs/cuda.h"
+
+// NORDC: define internal void @_Z4funcIiEvv()
+// NORDC: define void @_Z6kernelIiEvv()
+// RDC: define weak_odr void @_Z4funcIiEvv()
+// RDC: define weak_odr void @_Z6kernelIiEvv()
+
+template <typename T> __device__ void func() {}
+template <typename T> __global__ void kernel() {}
+
+template __device__ void func<int>();
+template __global__ void kernel<int>();
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -4483,13 +4483,16 @@
   // and must all be equivalent. However, we are not allowed to
   // throw away these explicit instantiations.
   //
-  // We don't currently support CUDA device code spread out across multiple TUs,
+  // CUDA/HIP: For -fno-gpu-rdc case, device code is limited to one TU,
   // so say that CUDA templates are either external (for kernels) or internal.
-  // This lets llvm perform aggressive inter-procedural optimizations.
+  // This lets llvm perform aggressive inter-procedural optimizations. For
+  // -fgpu-rdc case, device function calls across multiple TU's are allowed,
+  // therefore we need to follow the normal linkage paradigm.
   if (Linkage == GVA_StrongODR) {
-    if (Context.getLangOpts().AppleKext)
+    if (getLangOpts().AppleKext)
       return llvm::Function::ExternalLinkage;
-    if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice)
+    if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+        !getLangOpts().GPURelocatableDeviceCode)
       return D->hasAttr<CUDAGlobalAttr>() ? llvm::Function::ExternalLinkage
                                           : llvm::Function::InternalLinkage;
     return llvm::Function::WeakODRLinkage;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D90311.302548.patch
Type: text/x-patch
Size: 2268 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201103/d265081a/attachment.bin>