Skip to content

Commit

Permalink
Merged master:107c3a12d627 into amd-gfx:c198fccc05c7
Browse files Browse the repository at this point in the history 
Local branch amd-gfx c198fcc Merged master:e1af54296c52 into amd-gfx:1c7c56ddc669
Remote branch master 107c3a1 [WebAssembly] Implement ref.null
  • Loading branch information
Sw authored and Sw committed Nov 3, 2020
2 parents c198fcc + 107c3a1 commit 194fc97
Show file tree
Hide file tree
Showing 43 changed files with 16,367 additions and 1,184 deletions.
6 changes: 1 addition & 5 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8179,18 +8179,14 @@ def err_dynamic_var_init : Error<
"__device__, __constant__, and __shared__ variables.">;
def err_shared_var_init : Error<
"initialization is not supported for __shared__ variables.">;
def err_device_static_local_var : Error<
"within a %select{__device__|__global__|__host__|__host__ __device__}0 "
"function, only __shared__ variables or const variables without device "
"memory qualifier may be marked 'static'">;
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
def err_cuda_host_shared : Error<
"__shared__ local variables not allowed in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">;
def err_cuda_ovl_target : Error<
"%select{__device__|__global__|__host__|__host__ __device__}0 function %1 "
"cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
Expand Down
106 changes: 106 additions & 0 deletions clang/lib/Headers/__clang_hip_cmath.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,8 @@

#if defined(__cplusplus)
#include <limits>
#include <type_traits>
#include <utility>
#endif
#include <limits.h>
#include <stdint.h>
Expand Down Expand Up @@ -205,6 +207,72 @@ template <bool __B, class __T = void> struct __hip_enable_if {};

template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };

// decltype is only available in C++11 and above.
#if __cplusplus >= 201103L
// __hip_promote
namespace __hip {

template <class _Tp> struct __numeric_type {
static void __test(...);
static _Float16 __test(_Float16);
static float __test(float);
static double __test(char);
static double __test(int);
static double __test(unsigned);
static double __test(long);
static double __test(unsigned long);
static double __test(long long);
static double __test(unsigned long long);
static double __test(double);

typedef decltype(__test(std::declval<_Tp>())) type;
static const bool value = !std::is_same<type, void>::value;
};

template <> struct __numeric_type<void> { static const bool value = true; };

template <class _A1, class _A2 = void, class _A3 = void,
bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
&&__numeric_type<_A3>::value>
class __promote_imp {
public:
static const bool value = false;
};

template <class _A1, class _A2, class _A3>
class __promote_imp<_A1, _A2, _A3, true> {
private:
typedef typename __promote_imp<_A1>::type __type1;
typedef typename __promote_imp<_A2>::type __type2;
typedef typename __promote_imp<_A3>::type __type3;

public:
typedef decltype(__type1() + __type2() + __type3()) type;
static const bool value = true;
};

template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
private:
typedef typename __promote_imp<_A1>::type __type1;
typedef typename __promote_imp<_A2>::type __type2;

public:
typedef decltype(__type1() + __type2()) type;
static const bool value = true;
};

template <class _A1> class __promote_imp<_A1, void, void, true> {
public:
typedef typename __numeric_type<_A1>::type type;
static const bool value = true;
};

template <class _A1, class _A2 = void, class _A3 = void>
class __promote : public __promote_imp<_A1, _A2, _A3> {};

} // namespace __hip
#endif //__cplusplus >= 201103L

// __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
// avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
// floor(double).
Expand All @@ -219,6 +287,18 @@ template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
// __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
// or integer argument to avoid compilation error due to ambibuity. e.g.
// max(5.0f, 6.0) is resolved with max(double, double).
#if __cplusplus >= 201103L
#define __HIP_OVERLOAD2(__retty, __fn) \
template <typename __T1, typename __T2> \
__DEVICE__ typename __hip_enable_if< \
std::numeric_limits<__T1>::is_specialized && \
std::numeric_limits<__T2>::is_specialized, \
typename __hip::__promote<__T1, __T2>::type>::type \
__fn(__T1 __x, __T2 __y) { \
typedef typename __hip::__promote<__T1, __T2>::type __result_type; \
return __fn((__result_type)__x, (__result_type)__y); \
}
#else
#define __HIP_OVERLOAD2(__retty, __fn) \
template <typename __T1, typename __T2> \
__DEVICE__ \
Expand All @@ -228,6 +308,7 @@ template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
__fn(__T1 __x, __T2 __y) { \
return __fn((double)__x, (double)__y); \
}
#endif

__HIP_OVERLOAD1(double, abs)
__HIP_OVERLOAD1(double, acos)
Expand Down Expand Up @@ -296,6 +377,18 @@ __HIP_OVERLOAD2(double, max)
__HIP_OVERLOAD2(double, min)

// Additional Overloads that don't quite match HIP_OVERLOAD.
#if __cplusplus >= 201103L
template <typename __T1, typename __T2, typename __T3>
__DEVICE__ typename __hip_enable_if<
std::numeric_limits<__T1>::is_specialized &&
std::numeric_limits<__T2>::is_specialized &&
std::numeric_limits<__T3>::is_specialized,
typename __hip::__promote<__T1, __T2, __T3>::type>::type
fma(__T1 __x, __T2 __y, __T3 __z) {
typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
}
#else
template <typename __T1, typename __T2, typename __T3>
__DEVICE__
typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized &&
Expand All @@ -305,6 +398,7 @@ __DEVICE__
fma(__T1 __x, __T2 __y, __T3 __z) {
return ::fma((double)__x, (double)__y, (double)__z);
}
#endif

template <typename __T>
__DEVICE__
Expand All @@ -327,6 +421,17 @@ __DEVICE__
return ::modf((double)__x, __exp);
}

#if __cplusplus >= 201103L
template <typename __T1, typename __T2>
__DEVICE__
typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized &&
std::numeric_limits<__T2>::is_specialized,
typename __hip::__promote<__T1, __T2>::type>::type
remquo(__T1 __x, __T2 __y, int *__quo) {
typedef typename __hip::__promote<__T1, __T2>::type __result_type;
return ::remquo((__result_type)__x, (__result_type)__y, __quo);
}
#else
template <typename __T1, typename __T2>
__DEVICE__
typename __hip_enable_if<std::numeric_limits<__T1>::is_specialized &&
Expand All @@ -335,6 +440,7 @@ __DEVICE__
remquo(__T1 __x, __T2 __y, int *__quo) {
return ::remquo((double)__x, (double)__y, __quo);
}
#endif

template <typename __T>
__DEVICE__
Expand Down
1 change: 0 additions & 1 deletion clang/lib/Sema/SemaCUDA.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -519,7 +519,6 @@ void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
VD->hasAttr<CUDASharedAttr>()) {
if (LangOpts.GPUAllowDeviceInit)
return;
assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
bool AllowedInit = false;
if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
AllowedInit =
Expand Down
25 changes: 1 addition & 24 deletions clang/lib/Sema/SemaDecl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13172,32 +13172,9 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) {
}
}

if (VD->isStaticLocal()) {
if (VD->isStaticLocal())
CheckStaticLocalForDllExport(VD);

if (dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
// CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__
// function, only __shared__ variables or variables without any device
// memory qualifiers may be declared with static storage class.
// Note: It is unclear how a function-scope non-const static variable
// without device memory qualifier is implemented, therefore only static
// const variable without device memory qualifier is allowed.
[&]() {
if (!getLangOpts().CUDA)
return;
if (VD->hasAttr<CUDASharedAttr>())
return;
if (VD->getType().isConstQualified() &&
!(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
return;
if (CUDADiagIfDeviceCode(VD->getLocation(),
diag::err_device_static_local_var)
<< CurrentCUDATarget())
VD->setInvalidDecl();
}();
}
}

// Perform check for initializers of device-side global variables.
// CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
// 7.5). We must also apply the same checks to all __shared__
Expand Down
21 changes: 17 additions & 4 deletions clang/lib/Sema/SemaDeclAttr.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4394,8 +4394,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL))
return;
const auto *VD = cast<VarDecl>(D);
if (!VD->hasGlobalStorage()) {
S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant);
if (VD->hasLocalStorage()) {
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
Expand Down Expand Up @@ -4456,6 +4456,20 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
}

static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
if (checkAttrMutualExclusion<CUDAGlobalAttr>(S, D, AL)) {
return;
}

if (const auto *VD = dyn_cast<VarDecl>(D)) {
if (VD->hasLocalStorage()) {
S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
return;
}
}
D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL));
}

static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *Fn = cast<FunctionDecl>(D);
if (!Fn->isInlineSpecified()) {
Expand Down Expand Up @@ -7526,8 +7540,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleGlobalAttr(S, D, AL);
break;
case ParsedAttr::AT_CUDADevice:
handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
AL);
handleDeviceAttr(S, D, AL);
break;
case ParsedAttr::AT_CUDAHost:
handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
Expand Down
6 changes: 6 additions & 0 deletions clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,8 @@

// Test function scope static device variable, which should not be externalized.
// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
// DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43

// Check a static device variable referenced by host function is externalized.
// DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
Expand Down Expand Up @@ -78,6 +80,8 @@ inline __device__ void devfun(const int ** b) {

__global__ void kernel(int *a, const int **b) {
const static int w = 1;
const static __constant__ int local_static_constant = 42;
const static __device__ int local_static_device = 43;
a[0] = x;
a[1] = y;
a[2] = x2;
Expand All @@ -86,6 +90,8 @@ __global__ void kernel(int *a, const int **b) {
a[5] = x5;
b[0] = &w;
b[1] = &z2;
b[2] = &local_static_constant;
b[3] = &local_static_device;
devfun(b);
}

Expand Down
4 changes: 2 additions & 2 deletions clang/test/SemaCUDA/bad-attributes.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -64,11 +64,11 @@ __global__ static inline void foobar() {};

__constant__ int global_constant;
void host_fn() {
__constant__ int c; // expected-error {{__constant__ variables must be global}}
__constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
__shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
}
__device__ void device_fn() {
__constant__ int c; // expected-error {{__constant__ variables must be global}}
__constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
}

typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}
Expand Down
Loading

0 comments on commit 194fc97

Please sign in to comment.