diff options
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/lib/Basic/Targets/AMDGPU.cpp | 11 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/types.cu | 10 |
2 files changed, 21 insertions, 0 deletions
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 19dc05f5b0a..b04d894855e 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -307,5 +307,16 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts, } void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) { + assert(HalfFormat == Aux->HalfFormat); + assert(FloatFormat == Aux->FloatFormat); + assert(DoubleFormat == Aux->DoubleFormat); + + // On x86_64 long double is 80-bit extended precision format, which is + // not supported by AMDGPU. 128-bit floating point format is also not + // supported by AMDGPU. Therefore keep its own format for these two types. + auto SaveLongDoubleFormat = LongDoubleFormat; + auto SaveFloat128Format = Float128Format; copyAuxTarget(Aux); + LongDoubleFormat = SaveLongDoubleFormat; + Float128Format = SaveFloat128Format; } diff --git a/clang/test/CodeGenCUDA/types.cu b/clang/test/CodeGenCUDA/types.cu new file mode 100644 index 00000000000..43c2b85ea78 --- /dev/null +++ b/clang/test/CodeGenCUDA/types.cu @@ -0,0 +1,10 @@ +// RUN: %clang_cc1 -triple amdgcn -aux-triple x86_64 -fcuda-is-device -emit-llvm %s -o - | FileCheck -check-prefix=DEV %s +// RUN: %clang_cc1 -triple x86_64 -aux-triple amdgcn -emit-llvm %s -o - | FileCheck -check-prefix=HOST %s + +#include "Inputs/cuda.h" + +// HOST: @ld_host = global x86_fp80 0xK00000000000000000000 +long double ld_host; + +// DEV: @ld_device = addrspace(1) externally_initialized global double 0.000000e+00 +__device__ long double ld_device; |
