Index: lib/CodeGen/CGCUDABuiltin.cpp
===================================================================
--- lib/CodeGen/CGCUDABuiltin.cpp
+++ lib/CodeGen/CGCUDABuiltin.cpp
@@ -83,6 +83,11 @@
                E->arguments(), E->getDirectCallee(),
                /* ParamsToSkip = */ 0);
 
+  // We don't know how to emit non-scalar varargs, so just remove them.
+  Args.erase(std::remove_if(Args.begin() + 1, Args.end(),
+                            [](const CallArg &A) { return !A.RV.isScalar(); }),
+             Args.end());
+
   // Construct and fill the args buffer that we'll pass to vprintf.
   llvm::Value *BufferPtr;
   if (Args.size() <= 1) {
Index: test/CodeGenCUDA/printf.cu
===================================================================
--- test/CodeGenCUDA/printf.cu
+++ test/CodeGenCUDA/printf.cu
@@ -41,3 +41,12 @@
     printf("%d", 42);
   }
 }
+
+// Check that we don't crash when asked to printf a non-scalar arg.
+struct Struct {
+  int x;
+  int y;
+};
+__device__ void PrintfNonScalar() {
+  printf("%d", Struct());
+}