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

[Justin Lebar via cfe-commits]

This revision was automatically updated to reflect the committed changes.
jlebar marked an inline comment as done.
Closed by commit rL260479: [CUDA] Don't crash when trying to printf a 
non-scalar object. (authored by jlebar).

Changed prior to commit:
  http://reviews.llvm.org/D17103?vs=47569&id=47572#toc

Repository:
  rL LLVM

http://reviews.llvm.org/D17103

Files:
  cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
  cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu

Index: cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
+++ cfe/trunk/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::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) {
Index: cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
+++ cfe/trunk/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());
+}



Index: cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGCUDABuiltin.cpp
+++ cfe/trunk/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::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) {
Index: cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/printf-aggregate.cu
+++ cfe/trunk/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());
+}

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits