[PATCH] D17103: [CUDA] Don't crash when trying to printf a non-scalar object.

[Justin Lebar via cfe-commits]

jlebar created this revision.
jlebar added reviewers: majnemer, rnk.
jlebar added subscribers: tra, jhen, cfe-commits.

We can't do the right thing, since there's no right thing to do, but at
least we can not crash the compiler.

http://reviews.llvm.org/D17103

Files:
  lib/CodeGen/CGCUDABuiltin.cpp
  test/CodeGenCUDA/printf.cu

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


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17103.47533.patch
Type: text/x-patch
Size: 1047 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160210/fa9c0815/attachment.bin>