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

[Henry Linjamäki via Phabricator via cfe-commits]

linjamaki updated this revision to Diff 382228.
linjamaki added a comment.

Rephrase expressions.


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
@@ -1496,8 +1496,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
@@ -1855,7 +1855,8 @@
 }
 
 void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
-  if (!LangOpts.SYCLIsDevice && !(LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+  if (!(LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+        LangOpts.CUDAIsDevice))
     return;
 
   if (isUnevaluatedContext() || Ty.isNull())


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D111047.382228.patch
Type: text/x-patch
Size: 2539 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20211026/72a9be48/attachment.bin>