[PATCH] D111047: CUDA/HIP: Allow __int128 on the host side

[Artem Belevich via Phabricator via cfe-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rGc99b2c63169d: CUDA/HIP: Allow __int128 on the host side (authored by linjamaki, committed by tra).

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D111047

Files:
  clang/lib/Sema/Sema.cpp
  clang/lib/Sema/SemaType.cpp
  clang/test/SemaCUDA/allow-int128.cu
  clang/test/SemaCUDA/spirv-int128.cu


Index: clang/test/SemaCUDA/spirv-int128.cu
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/test/SemaCUDA/allow-int128.cu
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -1495,8 +1495,8 @@
   }
   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)
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -1941,7 +1941,8 @@
   };
 
   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();


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D111047.397420.patch
Type: text/x-patch
Size: 2562 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220105/974b7cac/attachment.bin>