Index: lib/CodeGen/CGCUDABuiltin.cpp =================================================================== --- lib/CodeGen/CGCUDABuiltin.cpp +++ lib/CodeGen/CGCUDABuiltin.cpp @@ -83,6 +83,13 @@ 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::getIgnored(); + } + // Construct and fill the args buffer that we'll pass to vprintf. llvm::Value *BufferPtr; if (Args.size() <= 1) { Index: test/CodeGenCUDA/printf-aggregate.cu =================================================================== --- /dev/null +++ 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()); +}