r282646 - [CUDA] Disallow exceptions in device code.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 28 15:45:54 PDT 2016
Author: jlebar
Date: Wed Sep 28 17:45:54 2016
New Revision: 282646
URL: http://llvm.org/viewvc/llvm-project?rev=282646&view=rev
Log:
[CUDA] Disallow exceptions in device code.
Reviewers: tra
Subscribers: cfe-commits, jhen
Differential Revision: https://reviews.llvm.org/D25036
Added:
cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
cfe/trunk/test/SemaCUDA/exceptions.cu
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaExprCXX.cpp
cfe/trunk/lib/Sema/SemaStmt.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=282646&r1=282645&r2=282646&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Wed Sep 28 17:45:54 2016
@@ -6702,6 +6702,9 @@ def err_cuda_unattributed_constexpr_cann
"attribute, or build with -fno-cuda-host-device-constexpr.">;
def note_cuda_conflicting_device_function_declared_here : Note<
"conflicting __device__ function declared here">;
+def err_cuda_device_exceptions : Error<
+ "cannot use '%0' in "
+ "%select{__device__|__global__|__host__|__host__ __device__}1 function %2">;
def err_dynamic_var_init : Error<
"dynamic initialization is not supported for "
"__device__, __constant__, and __shared__ variables.">;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=282646&r1=282645&r2=282646&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Sep 28 17:45:54 2016
@@ -9245,6 +9245,16 @@ public:
/// Otherwise, returns true without emitting any diagnostics.
bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
+ /// Check whether a 'try' or 'throw' expression is allowed within the current
+ /// context, and raise an error or create a deferred error, as appropriate.
+ ///
+ /// 'try' and 'throw' are never allowed in CUDA __device__ functions, and are
+ /// allowed in __host__ __device__ functions only if those functions are never
+ /// codegen'ed for the device.
+ ///
+ /// ExprTy should be the string "try" or "throw", as appropriate.
+ bool CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy);
+
/// Finds a function in \p Matches with highest calling priority
/// from \p Caller context and erases all functions with lower
/// calling priority.
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=282646&r1=282645&r2=282646&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Sep 28 17:45:54 2016
@@ -515,3 +515,27 @@ bool Sema::CheckCUDACall(SourceLocation
}
return true;
}
+
+bool Sema::CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy) {
+ assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+ FunctionDecl *CurFn = dyn_cast<FunctionDecl>(CurContext);
+ if (!CurFn)
+ return true;
+ CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn);
+
+ // Raise an error immediately if this is a __global__ or __device__ function.
+ // If it's a __host__ __device__ function, enqueue a deferred error which will
+ // be emitted if the function is codegen'ed for device.
+ if (Target == CFT_Global || Target == CFT_Device) {
+ Diag(Loc, diag::err_cuda_device_exceptions) << ExprTy << Target << CurFn;
+ return false;
+ }
+ if (Target == CFT_HostDevice && getLangOpts().CUDAIsDevice) {
+ PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()};
+ ErrPD.Reset(diag::err_cuda_device_exceptions);
+ ErrPD << ExprTy << Target << CurFn;
+ CurFn->addDeferredDiag({Loc, std::move(ErrPD)});
+ return false;
+ }
+ return true;
+}
Modified: cfe/trunk/lib/Sema/SemaExprCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExprCXX.cpp?rev=282646&r1=282645&r2=282646&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExprCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExprCXX.cpp Wed Sep 28 17:45:54 2016
@@ -683,6 +683,10 @@ ExprResult Sema::BuildCXXThrow(SourceLoc
!getSourceManager().isInSystemHeader(OpLoc))
Diag(OpLoc, diag::err_exceptions_disabled) << "throw";
+ // Exceptions aren't allowed in CUDA device code.
+ if (getLangOpts().CUDA)
+ CheckCUDAExceptionExpr(OpLoc, "throw");
+
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw";
Modified: cfe/trunk/lib/Sema/SemaStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaStmt.cpp?rev=282646&r1=282645&r2=282646&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaStmt.cpp (original)
+++ cfe/trunk/lib/Sema/SemaStmt.cpp Wed Sep 28 17:45:54 2016
@@ -3644,6 +3644,10 @@ StmtResult Sema::ActOnCXXTryBlock(Source
!getSourceManager().isInSystemHeader(TryLoc))
Diag(TryLoc, diag::err_exceptions_disabled) << "try";
+ // Exceptions aren't allowed in CUDA device code.
+ if (getLangOpts().CUDA)
+ CheckCUDAExceptionExpr(TryLoc, "try");
+
if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope())
Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try";
Added: cfe/trunk/test/SemaCUDA/exceptions-host-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions-host-device.cu?rev=282646&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/exceptions-host-device.cu (added)
+++ cfe/trunk/test/SemaCUDA/exceptions-host-device.cu Wed Sep 28 17:45:54 2016
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null
+// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null
+
+#include "Inputs/cuda.h"
+
+// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
+// function if and only if it's codegen'ed for device.
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__host__ __device__ void hd1() {
+ throw NULL;
+ try {} catch(void*) {}
+#ifndef HOST
+ // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function 'hd1'}}
+ // expected-error at -3 {{cannot use 'try' in __host__ __device__ function 'hd1'}}
+#endif
+}
+
+// No error, never instantiated on device.
+inline __host__ __device__ void hd2() {
+ throw NULL;
+ try {} catch(void*) {}
+}
+void call_hd2() { hd2(); }
+
+// Error, instantiated on device.
+inline __host__ __device__ void hd3() {
+ throw NULL;
+ try {} catch(void*) {}
+#ifndef HOST
+ // expected-error at -3 {{cannot use 'throw' in __host__ __device__ function 'hd3'}}
+ // expected-error at -3 {{cannot use 'try' in __host__ __device__ function 'hd3'}}
+#endif
+}
+__device__ void call_hd3() { hd3(); }
Added: cfe/trunk/test/SemaCUDA/exceptions.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/exceptions.cu?rev=282646&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/exceptions.cu (added)
+++ cfe/trunk/test/SemaCUDA/exceptions.cu Wed Sep 28 17:45:54 2016
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+void host() {
+ throw NULL;
+ try {} catch(void*) {}
+}
+__device__ void device() {
+ throw NULL;
+ // expected-error at -1 {{cannot use 'throw' in __device__ function 'device'}}
+ try {} catch(void*) {}
+ // expected-error at -1 {{cannot use 'try' in __device__ function 'device'}}
+}
+__global__ void kernel() {
+ throw NULL;
+ // expected-error at -1 {{cannot use 'throw' in __global__ function 'kernel'}}
+ try {} catch(void*) {}
+ // expected-error at -1 {{cannot use 'try' in __global__ function 'kernel'}}
+}
More information about the cfe-commits
mailing list