[clang] f2a1331 - [CUDA][HIP] Do not mark extern shared var (#65990)

[via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2023-09-11T17:04:55-04:00
New Revision: f2a1331a01ff8ad19fed6bd407501791467ad061

URL: https://github.com/llvm/llvm-project/commit/f2a1331a01ff8ad19fed6bd407501791467ad061
DIFF: https://github.com/llvm/llvm-project/commit/f2a1331a01ff8ad19fed6bd407501791467ad061.diff

LOG: [CUDA][HIP] Do not mark extern shared var (#65990)

Fixes: https://github.com/llvm/llvm-project/issues/65806

Currently clang put extern shared var ODR-used by host device functions
in global var __clang_gpu_used_external. This behavior was due to
https://reviews.llvm.org/D123441. However, clang should not do that for
extern shared vars since their addresses are per warp, therefore cannot
be accessed by host code.

Added: 
    

Modified: 
    clang/include/clang/AST/ASTContext.h
    clang/lib/Sema/SemaExpr.cpp
    clang/test/CodeGenCUDA/host-used-extern.cu

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index ec6b579870abbdc..4ee32c76a95d8e3 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -1141,6 +1141,9 @@ class ASTContext : public RefCountedBase<ASTContext> {
   mutable TagDecl *MSGuidTagDecl = nullptr;
 
   /// Keep track of CUDA/HIP device-side variables ODR-used by host code.
+  /// This does not include extern shared variables used by device host
+  /// functions as addresses of shared variables are per warp, therefore
+  /// cannot be accessed by host code.
   llvm::DenseSet<const VarDecl *> CUDADeviceVarODRUsedByHost;
 
   /// Keep track of CUDA/HIP external kernels or device variables ODR-used by

diff  --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 3e9c3fad03918f0..92496b03ecabe54 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -19136,6 +19136,7 @@ MarkVarDeclODRUsed(ValueDecl *V, SourceLocation Loc, Sema &SemaRef,
                                : diag::note_cuda_host_var);
       }
     } else if (VarTarget == Sema::CVT_Device &&
+               !Var->hasAttr<CUDASharedAttr>() &&
                (UserTarget == Sema::CFT_Host ||
                 UserTarget == Sema::CFT_HostDevice)) {
       // Record a CUDA/HIP device side variable if it is ODR-used

diff  --git a/clang/test/CodeGenCUDA/host-used-extern.cu b/clang/test/CodeGenCUDA/host-used-extern.cu
index c7edabf5c582b6f..e8f8e12aad47d1c 100644
--- a/clang/test/CodeGenCUDA/host-used-extern.cu
+++ b/clang/test/CodeGenCUDA/host-used-extern.cu
@@ -9,6 +9,11 @@
 // RUN:   -std=c++11 -emit-llvm -o - -target-cpu gfx906 \
 // RUN:   | FileCheck -check-prefixes=NEG,NORDC %s
 
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -x hip %s \
+// RUN:   -fgpu-rdc -std=c++11 -emit-llvm -o - \
+// RUN:   | FileCheck -check-prefix=HOST-NEG %s
+
+
 #include "Inputs/cuda.h"
 
 // CHECK-LABEL: @__clang_gpu_used_external = internal {{.*}}global
@@ -21,10 +26,13 @@
 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel3v
 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @var2
 // NEG-NOT: @__clang_gpu_used_external = {{.*}} @var3
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @ext_shvar
+// NEG-NOT: @__clang_gpu_used_external = {{.*}} @shvar
 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel1v
 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel4v
 // NORDC-NOT: @__clang_gpu_used_external = {{.*}} @var1
-
+// HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @ext_shvar
+// HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @shvar
 __global__ void kernel1();
 
 // kernel2 is not marked as used since it is a definition.
@@ -49,3 +57,13 @@ void test() {
   void *p = (void*)kernel4;
   use(&var1);
 }
+
+__global__ void test_lambda_using_extern_shared() {
+  extern __shared__ int ext_shvar[];
+  __shared__ int shvar[10];
+  auto lambda = [&]() {
+    ext_shvar[0] = 1;
+    shvar[0] = 2;
+  };
+  lambda();
+}