[clang] c99b2c6 - CUDA/HIP: Allow __int128 on the host side

[Artem Belevich via cfe-commits]

Author: Henry Linjamäki
Date: 2022-01-04T16:09:26-08:00
New Revision: c99b2c63169d5aa6499143078790cb3eb87dee45

URL: https://github.com/llvm/llvm-project/commit/c99b2c63169d5aa6499143078790cb3eb87dee45
DIFF: https://github.com/llvm/llvm-project/commit/c99b2c63169d5aa6499143078790cb3eb87dee45.diff

LOG: CUDA/HIP: Allow __int128 on the host side

Consider case where `__int128` type is supported by the host target but
not by a device target (e.g. spirv*). Clang emits an error message for
unsupported type even if the device code does not use it. This patch
fixes this issue by emitting the error message when the device code
attempts to use the unsupported type.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D111047

Added: 
    clang/test/SemaCUDA/allow-int128.cu
    clang/test/SemaCUDA/spirv-int128.cu

Modified: 
    clang/lib/Sema/Sema.cpp
    clang/lib/Sema/SemaType.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index ba69400fdbbfc..60f37c17c3f18 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1941,7 +1941,8 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
   };
 
   auto CheckType = [&](QualType Ty, bool IsRetTy = false) {
-    if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+    if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+        LangOpts.CUDAIsDevice)
       CheckDeviceType(Ty);
 
     QualType UnqualTy = Ty.getCanonicalType().getUnqualifiedType();

diff  --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 0b3154e6bcb61..57825fe3d79b2 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -1495,8 +1495,8 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
   }
   case DeclSpec::TST_int128:
     if (!S.Context.getTargetInfo().hasInt128Type() &&
-        !S.getLangOpts().SYCLIsDevice &&
-        !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
+        !(S.getLangOpts().SYCLIsDevice || S.getLangOpts().CUDAIsDevice ||
+          (S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)))
       S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
         << "__int128";
     if (DS.getTypeSpecSign() == TypeSpecifierSign::Unsigned)

diff  --git a/clang/test/SemaCUDA/allow-int128.cu b/clang/test/SemaCUDA/allow-int128.cu
new file mode 100644
index 0000000000000..eb7b7e7f52862
--- /dev/null
+++ b/clang/test/SemaCUDA/allow-int128.cu
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu \
+// RUN:   -fcuda-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple nvptx \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu \
+// RUN:   -fcuda-is-device -verify -fsyntax-only %s
+
+// expected-no-diagnostics
+#define __device__ __attribute__((device))
+
+__int128 h_glb;
+__device__ __int128 d_unused;
+__device__ __int128 d_glb;
+__device__ __int128 bar() {
+  return d_glb;
+}

diff  --git a/clang/test/SemaCUDA/spirv-int128.cu b/clang/test/SemaCUDA/spirv-int128.cu
new file mode 100644
index 0000000000000..b2ff5ae5f6922
--- /dev/null
+++ b/clang/test/SemaCUDA/spirv-int128.cu
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple spirv64 -aux-triple x86_64-unknown-linux-gnu \
+// RUN:   -fcuda-is-device -verify -fsyntax-only %s
+
+#define __device__ __attribute__((device))
+
+__int128 h_glb;
+
+__device__ __int128 d_unused;
+
+// expected-note at +1 {{'d_glb' defined here}}
+__device__ __int128 d_glb;
+
+__device__ __int128 bar() {
+  // expected-error at +1 {{'d_glb' requires 128 bit size '__int128' type support, but target 'spirv64' does not support it}}
+  return d_glb;
+}