maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

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
This commit is contained in:
Henry Linjamäki 2022-01-04 16:00:36 -08:00 committed by Artem Belevich
parent 822448635e
commit c99b2c6316
4 changed files with 36 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/Sema/Sema.cpp
View File

@ -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();

4
clang/lib/Sema/SemaType.cpp
View File

@ -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)

16
clang/test/SemaCUDA/allow-int128.cu Normal file
View File

@ -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;
}

16
clang/test/SemaCUDA/spirv-int128.cu Normal file
View File

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