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()); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
