[CUDA][HIP] Exclude external variables from constant promotion.

[PatriosTheGreat]

Promoting constant to external variables includes them to PTX which then leads to nvlinker failure.
See changes at device-use-host-var test.
Befor this change those variables was included to PTX without definition.

[llvmbot]

@llvm/pr-subscribers-clang

Author: Levon Ter-Grigoryan (PatriosTheGreat)

Changes

Promoting constant to external variables includes them to PTX which then leads to nvlinker failure.
See changes at device-use-host-var test.
Befor this change those variables was included to PTX without definition.

---

Full diff: https://github.com/llvm/llvm-project/pull/73549.diff

2 Files Affected:

• (modified) clang/lib/Sema/SemaCUDA.cpp (+1)
• (modified) clang/test/CodeGenCUDA/device-use-host-var.cu (+16)

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 318174f7be8fa95..f9d72e571e7b98b 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -783,6 +783,7 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
       (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
       !IsDependentVar(VD) &&
       ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
+       VD->getStorageClass() != SC_Extern &&
        HasAllowedCUDADeviceStaticInitializer(*this, VD,
                                              CICK_DeviceOrConstant))) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
diff --git a/clang/test/CodeGenCUDA/device-use-host-var.cu b/clang/test/CodeGenCUDA/device-use-host-var.cu
index 64de57e41b4b9f5..807a485f4c14972 100644
--- a/clang/test/CodeGenCUDA/device-use-host-var.cu
+++ b/clang/test/CodeGenCUDA/device-use-host-var.cu
@@ -2,6 +2,8 @@
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
 // RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
+// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda \
+// RUN:   -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
 
 #include "Inputs/cuda.h"
 
@@ -104,3 +106,17 @@ void fun() {
   (void) b<double>;
   (void) var_host_only;
 }
+
+// NEG-NOT: external_func
+extern __global__ void external_func();
+// NEG-NOT: @external_dep
+extern void* const external_dep[] = {
+  (void*)(external_func)
+};
+// NEG-NOT: @external_arr
+extern void* const external_arr[] = {};
+
+void* host_fun() {
+  (void) external_dep;
+  (void) external_arr;
+}

[Artem-B]

There are no CHECK lines here.

[PatriosTheGreat]

This case is checked by flag "-implicit-check-not=external_" at the line 6, so we can check that external is not mentioned anywhere at the device code.

[Artem-B]

This array is nomiannly host-only entity and should not be emitted on GPU at all, IMO. In fact, nvcc errors out if we attempt to access it on the GPU: https://godbolt.org/z/G15zn35Wd

Whether it's extern or not should not matter. I think.

@yxsamliu Sam, WDYT? I suspect there is/was a reason we may have allowed const access on both sides.

[yxsamliu]

It seems nvcc allows non-array type const var to be used in device code but not array type const var

https://godbolt.org/z/xjvbjPK77

I don't see why we cannot use array type const var in device code if we are able to emit them on device side. There may be CUDA/HIP code already using this feature. Disabling it may cause regressions.

On the other hand, I think disallow extern const var in device code is reasonable, since we do not know how it is initialized.

[arsenm]

Sounds broken that the behavior would differ between array and non-array ?

[arsenm]

Is #75799 related?