yaxunl updated this revision to Diff 329632.
yaxunl marked an inline comment as done.
yaxunl added a comment.
Follow C++ about ODR-use of variables.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D98193/new/
https://reviews.llvm.org/D98193
Files:
clang/lib/Sema/SemaExpr.cpp
clang/test/CodeGenCUDA/device-use-host-var.cu
clang/test/SemaCUDA/device-use-host-var.cu
Index: clang/test/SemaCUDA/device-use-host-var.cu
===================================================================
--- clang/test/SemaCUDA/device-use-host-var.cu
+++ clang/test/SemaCUDA/device-use-host-var.cu
@@ -5,37 +5,96 @@
#include "Inputs/cuda.h"
-int global_host_var;
+struct A {
+ int x;
+ static int host_var;
+};
+
+int A::host_var;
+
+namespace X {
+ int host_var;
+}
+
+static int static_host_var;
+
__device__ int global_dev_var;
__constant__ int global_constant_var;
__shared__ int global_shared_var;
-constexpr int global_constexpr_var = 1;
+
+int global_host_var;
const int global_const_var = 1;
+constexpr int global_constexpr_var = 1;
+
+int global_host_array[2] = {1, 2};
+const int global_const_array[2] = {1, 2};
+constexpr int global_constexpr_array[2] = {1, 2};
+
+A global_host_struct_var{1};
+const A global_const_struct_var{1};
+constexpr A global_constexpr_struct_var{1};
template<typename F>
__global__ void kernel(F f) { f(); } // dev-note2 {{called by 'kernel<(lambda}}
__device__ void dev_fun(int *out) {
- int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
+ // Check access device variables are allowed.
int &ref_dev_var = global_dev_var;
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
- const int &ref_constexpr_var = global_constexpr_var;
- const int &ref_const_var = global_const_var;
-
- *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
+ *out = ref_dev_var;
+ *out = ref_constant_var;
+ *out = ref_shared_var;
*out = global_dev_var;
*out = global_constant_var;
*out = global_shared_var;
- *out = global_constexpr_var;
+
+ // Check access of non-const host variables are not allowed.
+ *out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
*out = global_const_var;
+ *out = global_constexpr_var;
+ global_host_var = 1; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
+ // Check reference of non-constexpr host variables are not allowed.
+ int &ref_host_var = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
+ const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}}
+ const int &ref_constexpr_var = global_constexpr_var;
*out = ref_host_var;
- *out = ref_dev_var;
- *out = ref_constant_var;
- *out = ref_shared_var;
*out = ref_constexpr_var;
*out = ref_const_var;
+
+ // Check access member of non-constexpr struct type host variable is not allowed.
+ *out = global_host_struct_var.x; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}
+ *out = global_const_struct_var.x; // dev-error {{reference to __host__ variable 'global_const_struct_var' in __device__ function}}
+ *out = global_constexpr_struct_var.x;
+ global_host_struct_var.x = 1; // dev-error {{reference to __host__ variable 'global_host_struct_var' in __device__ function}}
+
+ // Check address taking of non-constexpr host variables is not allowed.
+ int *p = &global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __device__ function}}
+ const int *cp = &global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __device__ function}}
+ const int *cp2 = &global_constexpr_var;
+
+ // Check access elements of non-constexpr host array is not allowed.
+ *out = global_host_array[1]; // dev-error {{reference to __host__ variable 'global_host_array' in __device__ function}}
+ *out = global_const_array[1]; // dev-error {{reference to __host__ variable 'global_const_array' in __device__ function}}
+ *out = global_constexpr_array[1];
+
+ // Check ODR-use of host variables in namespace is not allowed.
+ *out = X::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}
+
+ // Check ODR-use of static host varables in class or file scope is not allowed.
+ *out = A::host_var; // dev-error {{reference to __host__ variable 'host_var' in __device__ function}}
+ *out = static_host_var; // dev-error {{reference to __host__ variable 'static_host_var' in __device__ function}}
+
+ // Check function-scope static variable is allowed.
+ static int static_var;
+ *out = static_var;
+
+ // Check non-ODR use of host varirables are allowed.
+ *out = sizeof(global_host_var);
+ *out = sizeof(global_host_struct_var.x);
+ decltype(global_host_var) var1;
+ decltype(global_host_struct_var.x) var2;
}
__global__ void global_fun(int *out) {
@@ -44,7 +103,7 @@
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
- const int &ref_const_var = global_const_var;
+ const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __global__ function}}
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __global__ function}}
*out = global_dev_var;
@@ -67,7 +126,7 @@
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
- const int &ref_const_var = global_const_var;
+ const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
*out = global_dev_var;
@@ -114,7 +173,7 @@
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
- const int &ref_const_var = global_const_var;
+ const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
// dev-error@-1 {{capture host variable 'out' by reference in device or host device lambda function}}
@@ -140,7 +199,7 @@
int &ref_constant_var = global_constant_var;
int &ref_shared_var = global_shared_var;
const int &ref_constexpr_var = global_constexpr_var;
- const int &ref_const_var = global_const_var;
+ const int &ref_const_var = global_const_var; // dev-error {{reference to __host__ variable 'global_const_var' in __host__ __device__ function}}
*out = global_host_var; // dev-error {{reference to __host__ variable 'global_host_var' in __host__ __device__ function}}
*out = global_dev_var;
@@ -166,7 +225,7 @@
template <class, int = 1, int = 1>
struct __attribute__((device_builtin_texture_type)) texture {
static texture<int> ref;
- __device__ int c() {
+ __device__ void c() {
auto &x = ref;
}
};
@@ -174,7 +233,15 @@
template <class, int = 1, int = 1>
struct not_a_texture {
static not_a_texture<int> ref;
- __device__ int c() {
+ __device__ void c() {
auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
}
};
+
+template<>
+not_a_texture<int> not_a_texture<int>::ref;
+
+__device__ void test_not_a_texture() {
+ not_a_texture<int> inst;
+ inst.c(); // dev-note {{in instantiation of member function 'not_a_texture<int, 1, 1>::c' requested here}}
+}
Index: clang/test/CodeGenCUDA/device-use-host-var.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/device-use-host-var.cu
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -emit-llvm -o - %s \
+// RUN: |FileCheck %s
+
+#include "Inputs/cuda.h"
+
+struct A {
+ int x;
+};
+
+constexpr int constexpr_var = 1;
+constexpr A constexpr_struct{2};
+constexpr A constexpr_array[4] = {0, 0, 0, 3};
+constexpr char constexpr_str[] = "abcd";
+const int const_var = 4;
+
+// CHECK: @_ZL13constexpr_str.const = private unnamed_addr constant [5 x i8] c"abcd\00"
+// CHECK: @_ZL13constexpr_var = internal constant i32 1
+// CHECK: @_ZL16constexpr_struct = internal constant %struct.A { i32 2 }
+// CHECK: @_ZL15constexpr_array = internal constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 3 }]
+// CHECK-NOT: external
+
+// CHECK: store i32 1
+// CHECK: store i32 2
+// CHECK: store i32 3
+// CHECK: store i32 4
+// CHECK: load i8, i8* getelementptr inbounds ([5 x i8], [5 x i8]* @_ZL13constexpr_str.const, i64 0, i64 3)
+// CHECK: store i32* @_ZL13constexpr_var
+// CHECK: store i32* getelementptr inbounds (%struct.A, %struct.A* @_ZL16constexpr_struct, i32 0, i32 0)
+// CHECK: store i32* getelementptr inbounds ([4 x %struct.A], [4 x %struct.A]* @_ZL15constexpr_array, i64 0, i64 3, i32 0)
+__device__ void dev_fun(int *out, const int **out2) {
+ *out = constexpr_var;
+ *out = constexpr_struct.x;
+ *out = constexpr_array[3].x;
+ *out = const_var;
+ *out = constexpr_str[3];
+ *out2 = &constexpr_var;
+ *out2 = &constexpr_struct.x;
+ *out2 = &constexpr_array[3].x;
+}
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -354,24 +354,6 @@
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
- // CUDA/HIP: Diagnose invalid references of host global variables in device
- // functions. Reference of device global variables in host functions is
- // allowed through shadow variables therefore it is not diagnosed.
- if (LangOpts.CUDAIsDevice) {
- auto *FD = dyn_cast_or_null<FunctionDecl>(CurContext);
- auto Target = IdentifyCUDATarget(FD);
- if (FD && Target != CFT_Host) {
- const auto *VD = dyn_cast<VarDecl>(D);
- if (VD && VD->hasGlobalStorage() && !VD->hasAttr<CUDADeviceAttr>() &&
- !VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
- !VD->getType()->isCUDADeviceBuiltinSurfaceType() &&
- !VD->getType()->isCUDADeviceBuiltinTextureType() &&
- !VD->isConstexpr() && !VD->getType().isConstQualified())
- targetDiag(*Locs.begin(), diag::err_ref_bad_target)
- << /*host*/ 2 << /*variable*/ 1 << VD << Target;
- }
- }
-
if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
if (auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);
@@ -17059,6 +17041,21 @@
CaptureType, DeclRefType,
FunctionScopeIndexToStopAt);
+ // Diagnose ODR-use of host global variables in device functions. Reference
+ // of device global variables in host functions is allowed through shadow
+ // variables therefore it is not diagnosed.
+ if (SemaRef.LangOpts.CUDAIsDevice) {
+ auto *FD = dyn_cast_or_null<FunctionDecl>(SemaRef.CurContext);
+ auto Target = SemaRef.IdentifyCUDATarget(FD);
+ if (Var && Var->isFileVarDecl() && !Var->hasAttr<CUDADeviceAttr>() &&
+ !Var->hasAttr<CUDAConstantAttr>() && !Var->hasAttr<CUDASharedAttr>() &&
+ !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
+ !Var->getType()->isCUDADeviceBuiltinTextureType()) {
+ SemaRef.targetDiag(Loc, diag::err_ref_bad_target)
+ << /*host*/ 2 << /*variable*/ 1 << Var << Target;
+ }
+ }
+
Var->markUsed(SemaRef.Context);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
