CUDA/HIP: Allow __int128 on the host side

[linjamaki]

Summary

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.

[linjamaki]

Rebase.

[tra]

LGTM. This could be useful for other types.
E.g. long double is technically not supported by NVPTX. We current demote it to double, so host-side does not cause errors, but ideally it should be treated similar to what you propose for i128 here.

clang/lib/Sema/Sema.cpp

1857–1858

Nit:I'd rephrase it a bit to make it a bit easier to understand (less negations).

clang/lib/Sema/SemaType.cpp

1497–1499

Similarly, here I'd separate the check for i128 support by the target and the compilation mode check.

[linjamaki]

Rephrase expressions.

[linjamaki]

Thanks for the review. Updated the patch with the style suggestions.

[linjamaki]

Rebase.

[linjamaki]

Rebase.

[linjamaki]

This patch should be ready to land. @tra, could you please commit it to the LLVM for us? Thanks.

clang/lib/Sema/Sema.cpp

Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID,
 
   SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice
                                  ? CUDADiagIfDeviceCode(Loc, DiagID)
                                  : CUDADiagIfHostCode(Loc, DiagID);
   SetIsLastErrorImmediate(DB.isImmediate());
   return DB;
 }
 
 void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
   if (isUnevaluatedContext() || Ty.isNull())
     return;
 
   Decl *C = cast<Decl>(getCurLexicalContext());
 
   // Memcpy operations for structs containing a member with unsupported type
   // are ok, though.
   if (const auto *MD = dyn_cast<CXXMethodDecl>(C)) {
     if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) &&
if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
           D->setInvalidDecl();
       }
       if (D)
         targetDiag(D->getLocation(), diag::note_defined_here, FD) << 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();
     const TargetInfo &TI = Context.getTargetInfo();
     if (!TI.hasLongDoubleType() && UnqualTy == Context.LongDoubleTy) {
       PartialDiagnostic PD = PDiag(diag::err_target_unsupported_type);
       if (D)
         PD << D;

clang/lib/Sema/SemaType.cpp

if (DS.getTypeSpecSign() == TypeSpecifierSign::Unsigned)
       Result = Context.getCorrespondingUnsignedType(Result);
 
     if (DS.isTypeSpecSat())
       Result = Context.getCorrespondingSaturatedType(Result);
 
     break;
   }
   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)
       Result = Context.UnsignedInt128Ty;
     else
       Result = Context.Int128Ty;
     break;
   case DeclSpec::TST_float16:

clang/test/SemaCUDA/allow-int128.cu

+// 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;
+}

clang/test/SemaCUDA/spirv-int128.cu

+// 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@+1 {{'d_glb' defined here}}
+__device__ __int128 d_glb;
+
+__device__ __int128 bar() {
+  // expected-error@+1 {{'d_glb' requires 128 bit size '__int128' type support, but target 'spirv64' does not support it}}
+  return d_glb;
+}