[PATCH] D21421: [NVPTX] Improve lowering of byval args of device functions.

Jingyue Wu via llvm-commits llvm-commits at lists.llvm.org
Wed Jun 15 20:32:22 PDT 2016


jingyue added inline comments.

================
Comment at: lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp:235
@@ +234,3 @@
+    if (Arg.getType()->isPointerTy() && Arg.hasByValAttr())
+      handleByValParam(&Arg);
+  return true;
----------------
Byval parameters of device functions might not be in the `.param` space (http://docs.nvidia.com/cuda/parallel-thread-execution/#parameter-state-space)

```
void device_function(struct S x) {
  ...
}

__global__ void kernel() {
  struct S x;
  device_func(x);
  ...
```


http://reviews.llvm.org/D21421





More information about the llvm-commits mailing list