yaxunl updated this revision to Diff 402924.
yaxunl added a comment.
fix test
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D118153/new/
https://reviews.llvm.org/D118153
Files:
clang/lib/AST/ExprConstant.cpp
clang/test/CodeGenCUDA/const-var.cu
clang/test/SemaCUDA/const-var.cu
Index: clang/test/SemaCUDA/const-var.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/const-var.cu
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN: -fsyntax-only -verify
+
+#include "Inputs/cuda.h"
+
+// Test const var initialized with address of a const var.
+// Both are promoted to device side.
+
+namespace Test1 {
+const int a = 1;
+
+struct B {
+ static const int *const p;
+ __device__ static const int *const p2;
+};
+const int *const B::p = &a;
+__device__ const int *const B::p2 = &a;
+
+__device__ void f() {
+ int y = a;
+ const int *x = B::p;
+ const int *z = B::p2;
+}
+}
+
+// Test const var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+namespace Test2 {
+int a = 1;
+// expected-note@-1{{host variable declared here}}
+
+struct B {
+ static int *const p;
+};
+int *const B::p = &a;
+// expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}}
+
+__device__ void f() {
+ int y = a;
+ // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}}
+ const int *x = B::p;
+ // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}}
+}
+}
+
+// Test device var initialized with address of a non-const host var.
+
+namespace Test3 {
+int a = 1;
+
+struct B {
+ __device__ static int *const p;
+};
+__device__ int *const B::p = &a;
+// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+}
Index: clang/test/CodeGenCUDA/const-var.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/const-var.cu
@@ -0,0 +1,45 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
+// RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s
+
+// Negative tests.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s
+
+#include "Inputs/cuda.h"
+
+// Test const var initialized with address of a const var.
+// Both are promoted to device side.
+
+// DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1
+// DEV-DAG: @_ZN5Test11B1pE = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
+// HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1
+// HOST-DAG: @_ZN5Test11B1pE = constant i32* @_ZN5Test1L1aE
+namespace Test1 {
+const int a = 1;
+
+struct B {
+ static const int *const p;
+};
+const int *const B::p = &a;
+}
+
+// Test const var initialized with address of a non-cost var.
+// Neither is promoted to device side.
+
+// DEV-NEG-NOT: @_ZN5Test2L1aE
+// DEV-NEG-NOT: @_ZN5Test21B1pE
+// HOST-DAG: @_ZN5Test21aE = global i32 1
+// HOST-DAG: @_ZN5Test21B1pE = constant i32* @_ZN5Test21aE
+
+namespace Test2 {
+int a = 1;
+
+struct B {
+ static int *const p;
+};
+int *const B::p = &a;
+}
Index: clang/lib/AST/ExprConstant.cpp
===================================================================
--- clang/lib/AST/ExprConstant.cpp
+++ clang/lib/AST/ExprConstant.cpp
@@ -983,6 +983,8 @@
discardCleanups();
}
+ ASTContext &getCtx() const override { return Ctx; }
+
void setEvaluatingDecl(APValue::LValueBase Base, APValue &Value,
EvaluatingDeclKind EDK = EvaluatingDeclKind::Ctor) {
EvaluatingDecl = Base;
@@ -1116,8 +1118,6 @@
Expr::EvalStatus &getEvalStatus() const override { return EvalStatus; }
- ASTContext &getCtx() const override { return Ctx; }
-
// If we have a prior diagnostic, it will be noting that the expression
// isn't a constant expression. This diagnostic is more important,
// unless we require this evaluation to produce a constant expression.
@@ -2216,6 +2216,19 @@
if (!isForManglingOnly(Kind) && Var->hasAttr<DLLImportAttr>())
// FIXME: Diagnostic!
return false;
+
+ // In CUDA/HIP device compilation, only device side variables have
+ // constant addresses.
+ if (Info.getCtx().getLangOpts().CUDA &&
+ Info.getCtx().getLangOpts().CUDAIsDevice) {
+ if (!Var->hasAttr<CUDADeviceAttr>() &&
+ !Var->hasAttr<CUDAConstantAttr>() &&
+ !Var->hasAttr<CUDASharedAttr>() &&
+ !Var->hasAttr<HIPManagedAttr>() &&
+ !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
+ !Var->getType()->isCUDADeviceBuiltinTextureType())
+ return false;
+ }
}
if (const auto *FD = dyn_cast<const FunctionDecl>(BaseVD)) {
// __declspec(dllimport) must be handled very carefully:
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
