From 05d90e2b083c240c83a4a60cfc93f914b67c2c8b Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Thu, 11 Feb 2016 02:00:52 +0000 Subject: [PATCH] [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 --- lib/CodeGen/CGCUDABuiltin.cpp | 7 +++++++ test/CodeGenCUDA/printf-aggregate.cu | 17 +++++++++++++++++ 2 files changed, 24 insertions(+) create mode 100644 test/CodeGenCUDA/printf-aggregate.cu diff --git a/lib/CodeGen/CGCUDABuiltin.cpp b/lib/CodeGen/CGCUDABuiltin.cpp index 0ccba8982a..ea3b888635 100644 --- a/lib/CodeGen/CGCUDABuiltin.cpp +++ b/lib/CodeGen/CGCUDABuiltin.cpp @@ -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) { diff --git a/test/CodeGenCUDA/printf-aggregate.cu b/test/CodeGenCUDA/printf-aggregate.cu new file mode 100644 index 0000000000..2e703b81d0 --- /dev/null +++ b/test/CodeGenCUDA/printf-aggregate.cu @@ -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()); +} -- 2.40.0