linjamaki created this revision.
Herald added a subscriber: yaxunl.
linjamaki edited the summary of this revision.
linjamaki added a reviewer: rsmith.
linjamaki added subscribers: bader, Anastasia.
linjamaki published this revision for review.
Herald added a reviewer: jdoerfert.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.
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.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D111047
Files:
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaExpr.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@+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 device '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,7 +1496,7 @@
}
case DeclSpec::TST_int128:
if (!S.Context.getTargetInfo().hasInt128Type() &&
- !S.getLangOpts().SYCLIsDevice &&
+ !S.getLangOpts().SYCLIsDevice && !S.getLangOpts().CUDAIsDevice &&
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__int128";
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -366,7 +366,8 @@
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+ LangOpts.CUDAIsDevice) {
if (auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -9569,7 +9569,8 @@
}
}
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+ LangOpts.CUDAIsDevice)
checkDeviceDecl(NewFD, D.getBeginLoc());
if (!getLangOpts().CPlusPlus) {
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@+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 device '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,7 +1496,7 @@
}
case DeclSpec::TST_int128:
if (!S.Context.getTargetInfo().hasInt128Type() &&
- !S.getLangOpts().SYCLIsDevice &&
+ !S.getLangOpts().SYCLIsDevice && !S.getLangOpts().CUDAIsDevice &&
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__int128";
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -366,7 +366,8 @@
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+ LangOpts.CUDAIsDevice) {
if (auto *VD = dyn_cast<ValueDecl>(D))
checkDeviceDecl(VD, Loc);
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -9569,7 +9569,8 @@
}
}
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+ LangOpts.CUDAIsDevice)
checkDeviceDecl(NewFD, D.getBeginLoc());
if (!getLangOpts().CPlusPlus) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
