r260479 - [CUDA] Don't crash when trying to printf a non-scalar object.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 10 18:00:53 PST 2016
Author: jlebar
Date: Wed Feb 10 20:00:52 2016
New Revision: 260479
URL: http://llvm.org/viewvc/llvm-project?rev=260479&view=rev
Log:
[CUDA] Don't crash when trying to printf a non-scalar object.
Summary:
We can't do the right thing, since there's no right thing to do, but at
least we can not crash the compiler.
Reviewers: majnemer, rnk
Subscribers: cfe-commits, jhen, tra
Differential Revision: http://reviews.llvm.org/D17103
Added:
cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
Modified:
cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
Modified: cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp?rev=260479&r1=260478&r2=260479&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp Wed Feb 10 20:00:52 2016
@@ -83,6 +83,13 @@ CodeGenFunction::EmitCUDADevicePrintfCal
E->arguments(), E->getDirectCallee(),
/* ParamsToSkip = */ 0);
+ // We don't know how to emit non-scalar varargs.
+ if (std::any_of(Args.begin() + 1, Args.end(),
+ [](const CallArg &A) { return !A.RV.isScalar(); })) {
+ CGM.ErrorUnsupported(E, "non-scalar arg to printf");
+ return RValue::get(llvm::ConstantInt::get(IntTy, 0));
+ }
+
// Construct and fill the args buffer that we'll pass to vprintf.
llvm::Value *BufferPtr;
if (Args.size() <= 1) {
Added: cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu?rev=260479&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu Wed Feb 10 20:00:52 2016
@@ -0,0 +1,17 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: not %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \
+// RUN: -o - %s 2>&1 | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check that we don't crash when asked to printf a non-scalar arg.
+struct Struct {
+ int x;
+ int y;
+};
+__device__ void PrintfNonScalar() {
+ // CHECK: cannot compile this non-scalar arg to printf
+ printf("%d", Struct());
+}
More information about the cfe-commits
mailing list