[clang] d880557 - [CUDA][HIP] Allow non-ODR use of host var in device
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 19 11:45:42 PDT 2021
Author: Yaxun (Sam) Liu
Date: 2021-04-19T14:45:24-04:00
New Revision: d8805574c183484f055552855fa82d2e8932415e
URL: https://github.com/llvm/llvm-project/commit/d8805574c183484f055552855fa82d2e8932415e
DIFF: https://github.com/llvm/llvm-project/commit/d8805574c183484f055552855fa82d2e8932415e.diff
LOG: [CUDA][HIP] Allow non-ODR use of host var in device
Reviewed by: Artem Belevich, Richard Smith
Differential Revision: https://reviews.llvm.org/D98193
Added:
clang/test/CodeGenCUDA/device-use-host-var.cu
Modified:
clang/lib/Headers/__clang_hip_math.h
clang/lib/Sema/SemaExpr.cpp
clang/test/Headers/hip-header.hip
clang/test/SemaCUDA/device-use-host-var.cu
Removed:
################################################################################
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index 35cf0ad3ba6c5..1f0982d92eff3 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -38,7 +38,7 @@ template<bool>
struct __compare_result{};
template<>
struct __compare_result<true> {
- static const bool valid;
+ static const __device__ bool valid;
};
__DEVICE__
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 38c25ca65374f..4c44295e7e142 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -355,24 +355,6 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
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);
@@ -17143,6 +17125,31 @@ MarkVarDeclODRUsed(VarDecl *Var, SourceLocation Loc, Sema &SemaRef,
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.CUDA && SemaRef.LangOpts.CUDAIsDevice) {
+ auto *FD = dyn_cast_or_null<FunctionDecl>(SemaRef.CurContext);
+ auto Target = SemaRef.IdentifyCUDATarget(FD);
+ auto IsEmittedOnDeviceSide = [](VarDecl *Var) {
+ if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
+ Var->hasAttr<CUDASharedAttr>() ||
+ Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ Var->getType()->isCUDADeviceBuiltinTextureType())
+ return true;
+ // Function-scope static variable in device functions or kernels are
+ // emitted on device side.
+ if (auto *FD = dyn_cast<FunctionDecl>(Var->getDeclContext())) {
+ return FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>();
+ }
+ return false;
+ };
+ if (Var && Var->hasGlobalStorage() && !IsEmittedOnDeviceSide(Var)) {
+ SemaRef.targetDiag(Loc, diag::err_ref_bad_target)
+ << /*host*/ 2 << /*variable*/ 1 << Var << Target;
+ }
+ }
+
Var->markUsed(SemaRef.Context);
}
diff --git a/clang/test/CodeGenCUDA/device-use-host-var.cu b/clang/test/CodeGenCUDA/device-use-host-var.cu
new file mode 100644
index 0000000000000..40dcef89bf5bf
--- /dev/null
+++ b/clang/test/CodeGenCUDA/device-use-host-var.cu
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | 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-DAG: @_ZL13constexpr_str.const = private unnamed_addr addrspace(4) constant [5 x i8] c"abcd\00"
+// CHECK-DAG: @_ZL13constexpr_var = internal addrspace(4) constant i32 1
+// CHECK-DAG: @_ZL16constexpr_struct = internal addrspace(4) constant %struct.A { i32 2 }
+// CHECK-DAG: @_ZL15constexpr_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 3 }]
+// CHECK-NOT: external
+
+// CHECK-LABEL: define{{.*}}@_Z7dev_funPiPPKi
+// CHECK: store i32 1
+// CHECK: store i32 2
+// CHECK: store i32 3
+// CHECK: store i32 4
+// CHECK: load i8, i8* getelementptr {{.*}} @_ZL13constexpr_str.const
+// CHECK: store i32* {{.*}}@_ZL13constexpr_var
+// CHECK: store i32* getelementptr {{.*}} @_ZL16constexpr_struct
+// CHECK: store i32* getelementptr {{.*}} @_ZL15constexpr_array
+__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;
+}
diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip
index 943254b63f2ed..c0ca394c613ce 100644
--- a/clang/test/Headers/hip-header.hip
+++ b/clang/test/Headers/hip-header.hip
@@ -5,6 +5,13 @@
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck %s
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN: -internal-isystem %S/Inputs/include \
+// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
+// RUN: -D__HIPCC_RTC__ -std=c++14 | FileCheck %s
+
// expected-no-diagnostics
diff --git a/clang/test/SemaCUDA/device-use-host-var.cu b/clang/test/SemaCUDA/device-use-host-var.cu
index c8ef7dbbb18dc..6e48544787438 100644
--- a/clang/test/SemaCUDA/device-use-host-var.cu
+++ b/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 @@ __global__ void global_fun(int *out) {
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 @@ __host__ __device__ void host_dev_fun(int *out) {
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 @@ void dev_lambda_capture_by_ref(int *out) {
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 at -1 {{capture host variable 'out' by reference in device or host device lambda function}}
@@ -140,7 +199,7 @@ void dev_lambda_capture_by_copy(int *out) {
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 @@ void dev_lambda_capture_by_copy(int *out) {
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,40 @@ struct __attribute__((device_builtin_texture_type)) texture {
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}}
+}
+
+// Test static variable in host function used by device function.
+void test_static_var_host() {
+ for (int i = 0; i < 10; i++) {
+ static int x;
+ struct A {
+ __device__ int f() {
+ return x; // dev-error{{reference to __host__ variable 'x' in __device__ function}}
+ }
+ };
+ }
+}
+
+// Test static variable in device function used by device function.
+__device__ void test_static_var_device() {
+ for (int i = 0; i < 10; i++) {
+ static int x;
+ int y = x;
+ struct A {
+ __device__ int f() {
+ return x;
+ }
+ };
+ }
+}
More information about the cfe-commits
mailing list