Parameters to CUDA kernels

Question

When invoking a CUDA kernel for a specific thread configuration, are there any strict rules on which memory space (device/host) kernel parameters should reside in and what type they should be?

Suppose I launch a 1-D grid of threads with

kernel<<<numblocks, threadsperblock >>> (/*parameters*/)

Can I pass an integer parameter int foo which is a host-integer variable, directly to the CUDA kernel? Or should I cudaMalloc memory for a single integer say dev_foo and then cudaMemcpy foo into devfoo and then pass devfoo as a kernel parameter?

score 13 · Accepted Answer · edited Nov 29 '11 at 05:25

13

The rules for kernel arguments are a logical consequence of C++ parameter passing rules and the fact that device and host memory are physically separate.

CUDA does not allow passing arguments by reference and you must be careful with pointers.

Specifically, you must pass parameters by value. Passing user-defined types requires that the default copy-constructor or your own copy-constructor (if present) does not contain any memory allocations (heap allocations with "new" or "malloc").

In summary pass-by-value works well for integral, floating point or other primitive types, and simple flat user-defined structs or class objects.

edited Nov 29 '11 at 05:25

harrism

26,505
2
57
88

answered Nov 28 '11 at 22:51

Yappie

399
2
8

1

CUDA has come a long way since this answer was written. Even when the answer was written I believe it should have been legal/possible to use pass-by-reference if the underlying data was pinned and mapped. Today it is legal/possible to use pass-by-reference with [UM/managed memory](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-unified-memory-programming-hd) and also with pinned/mapped memory (e.g. `cudaHostAlloc()`) – Robert Crovella Mar 18 '18 at 22:20

score 5 · Answer 2 · answered Nov 28 '11 at 22:06

You only need to use cudaMalloc() and cudaMemcpy() for blocks of data. Not single ints and the like. You also can pass structs as parameters, as long as they have no members pointing to a block of data in host memory.

So as a rule of thumb: if you are passing a pointer to a kernel, make sure it points into device memory.

Parameters to CUDA kernels

2 Answers2

Linked

Related