Index: clang/lib/Basic/Targets/AMDGPU.cpp =================================================================== --- clang/lib/Basic/Targets/AMDGPU.cpp +++ clang/lib/Basic/Targets/AMDGPU.cpp @@ -462,9 +462,13 @@ // supported by AMDGPU. Therefore keep its own format for these two types. auto SaveLongDoubleFormat = LongDoubleFormat; auto SaveFloat128Format = Float128Format; + auto SaveLongDoubleWidth = LongDoubleWidth; + auto SaveLongDoubleAlign = LongDoubleAlign; copyAuxTarget(Aux); LongDoubleFormat = SaveLongDoubleFormat; Float128Format = SaveFloat128Format; + LongDoubleWidth = SaveLongDoubleWidth; + LongDoubleAlign = SaveLongDoubleAlign; // For certain builtin types support on the host target, claim they are // support to pass the compilation of the host code during the device-side // compilation. Index: clang/test/CodeGenCUDA/long-double.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/long-double.cu @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 \ +// RUN: -aux-triple x86_64-unknown-gnu-linux -fcuda-is-device \ +// RUN: -emit-llvm -o - -x hip %s 2>&1 | FileCheck %s + +// RUN: %clang_cc1 -triple nvptx \ +// RUN: -aux-triple x86_64-unknown-gnu-linux -fcuda-is-device \ +// RUN: -emit-llvm -o - %s 2>&1 | FileCheck %s + +// CHECK: @_ZN15infinity_helperIeE5valueE = {{.*}} double 0x47EFFFFFD586B834, align 8 +// CHECK: @size = {{.*}} i32 8 + +#include "Inputs/cuda.h" + +template struct infinity_helper {}; +template <> struct infinity_helper { static constexpr long double value = 3.4028234e38L; }; +constexpr long double infinity_helper::value; +__device__ int size = sizeof(long double);