yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added subscribers: kosarev, kerbowa, tpr, jvesely.
Herald added a project: All.
yaxunl requested review of this revision.
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.
https://reviews.llvm.org/D127771
Files:
clang/lib/Basic/Targets/AMDGPU.cpp
clang/test/CodeGenCUDA/long-double.cu
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 <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);
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 <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);
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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
