[clang] [CUDA][HIP] Exclude external variables from constant promotion. (PR #73549)
Levon Ter-Grigoryan via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 27 09:34:09 PST 2023
https://github.com/PatriosTheGreat created https://github.com/llvm/llvm-project/pull/73549
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.
>From a1a50db9bd14467ca4463ec05affc381ab0ea1aa Mon Sep 17 00:00:00 2001
From: Levon Ter-Grigoryan <leonidusspartus at gmail.com>
Date: Mon, 27 Nov 2023 18:09:22 +0100
Subject: [PATCH] [CUDA][HIP] Exclude external variables from constant
promotion.
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.
---
clang/lib/Sema/SemaCUDA.cpp | 1 +
clang/test/CodeGenCUDA/device-use-host-var.cu | 16 ++++++++++++++++
2 files changed, 17 insertions(+)
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;
+}
More information about the cfe-commits
mailing list