[HIP] fix long double size

[yaxunl]

Summary

For amdgpu target long double type is the same as double type.
The width and align of long double type was incorrectly
overridden when copying aux target properties, which
caused assertion in codegen when emitting global
variables with long double type.

This patch fix that by saving and restoring width
and align of long double type.

[tra]

AFAICT, the test case you've added works fine with the compiler at HEAD: https://cuda.godbolt.org/z/q3xYMfdeb
I guess it only shows up in assertion-enabled builds. Can you check what happens if you run the test case compiled as CUDA? I suspect it will have the same issue -- we don't override setAuxTarget at all.

[yaxunl]

In D127771#3582640, @tra wrote:

AFAICT, the test case you've added works fine with the compiler at HEAD: https://cuda.godbolt.org/z/q3xYMfdeb
I guess it only shows up in assertion-enabled builds. Can you check what happens if you run the test case compiled as CUDA? I suspect it will have the same issue -- we don't override setAuxTarget at all.

The assertion only happens with assertion-enabled builds. However, the test is able to detect the issue with assertion-disabled builds since it checks the sizeof(long double). In the old build it is 16.

The lit test contains run line for nvptx. It seems nvptx is not subject to this issue since sizeof(long double)==8.

clang/lib/Basic/Targets/AMDGPU.cpp

void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) {
   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;
+  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.
   // FIXME: As the side effect, we also accept `__float128` uses in the device
   // code. To rejct these builtin types supported in the host target but not in
   // the device target, one approach would support `device_builtin` attribute
   // so that we could tell the device builtin types from the host ones. The
   // also solves the different representations of the same builtin type, such
   // as `size_t` in the MSVC environment.
   if (Aux->hasFloat128Type()) {
     HasFloat128 = true;
     Float128Format = DoubleFormat;
   }
 }

clang/test/CodeGenCUDA/long-double.cu

+// 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 <class> struct infinity_helper {};
+template <> struct infinity_helper<long double> { static constexpr long double value = 3.4028234e38L; };
+constexpr long double infinity_helper<long double>::value;
+__device__ int size = sizeof(long double);