Author: yaxunl Date: Thu Jan 31 13:57:51 2019 New Revision: 352801 URL: http://llvm.org/viewvc/llvm-project?rev=352801&view=rev Log: Do not copy long double and 128-bit fp format from aux target for AMDGPU
rC352620 caused regressions because it copied floating point format from aux target. floating point format decides whether extended long double is supported. It is x86_fp80 on x86 but IEEE double on amdgcn. Document usage of long doubel type in HIP programming guide https://github.com/ROCm-Developer-Tools/HIP/pull/890 Differential Revision: https://reviews.llvm.org/D57527 Added: cfe/trunk/test/CodeGenCUDA/types.cu Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=352801&r1=352800&r2=352801&view=diff ============================================================================== --- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original) +++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Thu Jan 31 13:57:51 2019 @@ -307,5 +307,16 @@ void AMDGPUTargetInfo::getTargetDefines( } 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; } Added: cfe/trunk/test/CodeGenCUDA/types.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/types.cu?rev=352801&view=auto ============================================================================== --- cfe/trunk/test/CodeGenCUDA/types.cu (added) +++ cfe/trunk/test/CodeGenCUDA/types.cu Thu Jan 31 13:57:51 2019 @@ -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; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits