yaxunl updated this revision to Diff 403706.
yaxunl added a comment.
Fix the regression in lit tests.
Basically in device compilation we still evaluate constant expression for host
functions or host template instantiation. If we just disallow host variable in
any constant expressions we will get errors in template class instantiation
which use host variables as non-type template arguments.
Therefore we should only disallow host variables in constant expressions in
situations when we are sure that allowing them will lead to issues, e.g. when
promoting const variables.
A CUDAConstantEvaluationContext is introduced in ASTContext to control this.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D118153/new/
https://reviews.llvm.org/D118153
Files:
clang/include/clang/AST/ASTContext.h
clang/lib/AST/ExprConstant.cpp
clang/lib/Sema/SemaCUDA.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/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -590,6 +590,8 @@
};
auto IsConstantInit = [&](const Expr *Init) {
assert(Init);
+ ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.Context,
+ /*NoWronSidedVars=*/true);
return Init->isConstantInitializer(S.Context,
VD->getType()->isReferenceType());
};
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 &&
+ Info.getCtx().CUDAConstantEvalCtx.NoWrongSidedVars) {
+ if (!Var->hasAttr<CUDADeviceAttr>() &&
+ !Var->hasAttr<CUDAConstantAttr>() &&
+ !Var->hasAttr<CUDASharedAttr>() &&
+ !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
+ !Var->getType()->isCUDADeviceBuiltinTextureType())
+ return false;
+ }
}
if (const auto *FD = dyn_cast<const FunctionDecl>(BaseVD)) {
// __declspec(dllimport) must be handled very carefully:
Index: clang/include/clang/AST/ASTContext.h
===================================================================
--- clang/include/clang/AST/ASTContext.h
+++ clang/include/clang/AST/ASTContext.h
@@ -653,6 +653,20 @@
/// Returns the clang bytecode interpreter context.
interp::Context &getInterpContext();
+ struct CUDAConstantEvalContext {
+ /// Do not allow wrong-sided variables in constant expressions.
+ bool NoWrongSidedVars = false;
+ } CUDAConstantEvalCtx;
+ struct CUDAConstantEvalContextRAII {
+ ASTContext &Ctx;
+ CUDAConstantEvalContext SavedCtx;
+ CUDAConstantEvalContextRAII(ASTContext &Ctx_, bool NoWrongSidedVars)
+ : Ctx(Ctx_), SavedCtx(Ctx_.CUDAConstantEvalCtx) {
+ Ctx_.CUDAConstantEvalCtx.NoWrongSidedVars = NoWrongSidedVars;
+ }
+ ~CUDAConstantEvalContextRAII() { Ctx.CUDAConstantEvalCtx = SavedCtx; }
+ };
+
/// Returns the dynamic AST node parent map context.
ParentMapContext &getParentMapContext();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
