dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@260479 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Justin Lebar 2016-02-11 02:00:52 +00:00
parent bdcdc83f36
commit 05d90e2b08
2 changed files with 24 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
lib/CodeGen/CGCUDABuiltin.cpp
View File

@ -83,6 +83,13 @@ CodeGenFunction::EmitCUDADevicePrintfCallExpr(const CallExpr *E,
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) {

17
test/CodeGenCUDA/printf-aggregate.cu Normal file
View File

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