[PATCH] D57527: Do not copy long double and 128-bit fp format from aux target for AMDGPU

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 184569.
yaxunl retitled this revision from "Do not copy floating pointer format from 
aux target" to "Do not copy long double and 128-bit fp format from aux target 
for AMDGPU".
yaxunl edited the summary of this revision.
yaxunl added a comment.
Herald added subscribers: t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, 
kzhuravl.

Fix in AMDGPUTargetInfo.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D57527/new/

https://reviews.llvm.org/D57527

Files:
  lib/Basic/Targets/AMDGPU.cpp
  test/CodeGenCUDA/types.cu


Index: test/CodeGenCUDA/types.cu
===================================================================
--- /dev/null
+++ 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;
Index: lib/Basic/Targets/AMDGPU.cpp
===================================================================
--- lib/Basic/Targets/AMDGPU.cpp
+++ lib/Basic/Targets/AMDGPU.cpp
@@ -307,5 +307,12 @@
 }
 
 void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) {
+  assert(HalfFormat == Aux->HalfFormat);
+  assert(FloatFormat == Aux->FloatFormat);
+  assert(DoubleFormat == Aux->DoubleFormat);
+  auto SaveLongDoubleFormat = LongDoubleFormat;
+  auto SaveFloat128Format = Float128Format;
   copyAuxTarget(Aux);
+  LongDoubleFormat = SaveLongDoubleFormat;
+  Float128Format = SaveFloat128Format;
 }



Index: test/CodeGenCUDA/types.cu
===================================================================
--- /dev/null
+++ 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;
Index: lib/Basic/Targets/AMDGPU.cpp
===================================================================
--- lib/Basic/Targets/AMDGPU.cpp
+++ lib/Basic/Targets/AMDGPU.cpp
@@ -307,5 +307,12 @@
 }
 
 void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) {
+  assert(HalfFormat == Aux->HalfFormat);
+  assert(FloatFormat == Aux->FloatFormat);
+  assert(DoubleFormat == Aux->DoubleFormat);
+  auto SaveLongDoubleFormat = LongDoubleFormat;
+  auto SaveFloat128Format = Float128Format;
   copyAuxTarget(Aux);
+  LongDoubleFormat = SaveLongDoubleFormat;
+  Float128Format = SaveFloat128Format;
 }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits