[clang] [CUDA][HIP] Exclude external variables from constant promotion. (PR #73549)

Yaxun Liu via cfe-commits cfe-commits at lists.llvm.org
Thu Dec 7 08:16:56 PST 2023


================
@@ -104,3 +106,14 @@ void fun() {
   (void) b<double>;
   (void) var_host_only;
 }
+
+extern __global__ void external_func();
+extern void* const external_dep[] = {
----------------
yxsamliu wrote:

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.

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


More information about the cfe-commits mailing list