reference argument on Cuda
cuda, nvidia, parallel-processing
Solution
A reference argument on a `__global__` function would not work, because passing the argument by reference in effect creates a pointer that the function will use to reference the argument. However this would usually result in dereferencing a host pointer on the device code, which is not allowed.
A `__device__` function may use reference arguments, however, because dereferencing a device pointer in device code is legal.
Regarding a "solution", just pass a pointer:
ExtractDisparityKernel ( ExtractDisparity *es)
And of course, make sure the argument you pass is a proper `cudaMalloc`-created pointer.
Generally speaking, contrary to the comment below, it is possible to use reference arguments on child kernels in a CDP (CUDA Dynamic Parallelism) setting. Certain types of reference arguments may still be unusable due to other limitations in a CDP setting, such as the limitation on use of the local memory space of the parent kernel.
And with the advent of unified memory (UM), it is possible now to use reference arguments even on kernels invoked from host code, with proper use of UM. Similarly, reference arguments referring to pinned host memory can be used.
Problem
I didn't know cuda doesn't support reference arguments. There are these two functions in my program: ``` __global__ void ExtractDisparityKernel ( ExtractDisparity& es) { es (); } __device__ __forceinline__ void computeAdjacentValue (int x1, int y1, int x2, int y2, float& value ) { .... } ``` Given the global function, the compiler reports error: /home/lv/pcl-trunk/gpu/kinfu_large_scale/src/cuda/estimate_combined.cu(959): error: a global routine cannot have reference arguments I searched for some solutions. Somebody says it is not allowed. But the device function doesn't report such kind of errors. I'm confused that whether cuda supports reference argument. Or the compiler is fooled somehow. Can anyone give a complete answer to this problem: where reference is allowed and not allowed?