Do not copy long double and 128-bit fp format from aux target for AMDGPU

[yaxunl]

Summary

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

[rjmccall]

Okay, so you silently have an incompatible ABI for anything in the system headers that mentions long double. Do you have any plans to address or work around that, or is the hope that it just doesn't matter?

I feel like this should be a special case for AMDGPU rather than a general behavior with aux targets.

[yaxunl]

In D57527#1379065, @rjmccall wrote:

Okay, so you silently have an incompatible ABI for anything in the system headers that mentions long double. Do you have any plans to address or work around that, or is the hope that it just doesn't matter?

I feel like this should be a special case for AMDGPU rather than a general behavior with aux targets.

If host do not pass long double to device we will be fine. So we need to diagnose long double kernel arguments. However I'd like to do it in separate patch since we want to fix the regression first.

Since this maybe a special case for AMDGPU, I will fix it in AMDGPUTargetInfo.

[rjmccall]

In D57527#1379088, @yaxunl wrote:

In D57527#1379065, @rjmccall wrote:

Okay, so you silently have an incompatible ABI for anything in the system headers that mentions long double. Do you have any plans to address or work around that, or is the hope that it just doesn't matter?

I feel like this should be a special case for AMDGPU rather than a general behavior with aux targets.

If host do not pass long double to device we will be fine. So we need to diagnose long double kernel arguments. However I'd like to do it in separate patch since we want to fix the regression first.

Okay. Do you also need to look for global structs and other way that information might be passed? I suppose at some level you just have to document it as a danger and treat further diagnostics as QoI.

Since this maybe a special case for AMDGPU, I will fix it in AMDGPUTargetInfo.

Alright. That should be as easy as saving the old value and restoring it after the overwrite.

[yaxunl]

Fix in AMDGPUTargetInfo.

[yaxunl]

In D57527#1379159, @rjmccall wrote:

In D57527#1379088, @yaxunl wrote:

In D57527#1379065, @rjmccall wrote:

Okay, so you silently have an incompatible ABI for anything in the system headers that mentions long double. Do you have any plans to address or work around that, or is the hope that it just doesn't matter?

I feel like this should be a special case for AMDGPU rather than a general behavior with aux targets.

If host do not pass long double to device we will be fine. So we need to diagnose long double kernel arguments. However I'd like to do it in separate patch since we want to fix the regression first.

Okay. Do you also need to look for global structs and other way that information might be passed? I suppose at some level you just have to document it as a danger and treat further diagnostics as QoI.

I created a pull request to document long double usage in HIP https://github.com/ROCm-Developer-Tools/HIP/pull/890

[rjmccall]

Explanatory comment, please. Otherwise LGTM.

[yaxunl]

In D57527#1379287, @rjmccall wrote:

Explanatory comment, please. Otherwise LGTM.

will do when committing.

lib/Basic/Targets/AMDGPU.cpp

if (hasLDEXPF())
     Builder.defineMacro("__HAS_LDEXPF__");
   if (hasFP64())
     Builder.defineMacro("__HAS_FP64__");
   if (hasFastFMA())
     Builder.defineMacro("FP_FAST_FMA");
 }
 
 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;
 }

test/CodeGenCUDA/types.cu

+// 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;