CODEWITHSUNDEEP
cuda
gct
gct

Reputation: 14573

Passing value from device memory as kernel parameter in CUDA

I'm writing a CUDA application that has a step where the variance of some complex-valued input data is computed, and then that variance is used to threshold the data. I've got a reduction kernel that computes the variance for me, but I'm not sure if I have to pull the value back to the host to pass it to the thresholding kernel or not.

Is there a way to pass the value directly from device memory?

Upvotes: 0

Views: 1220

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 151879

You can use a __device__ variable to hold the variance value in-between kernel calls.

Put this before the definition of the kernels that use it:

__device__ float my_variance = 0.0f;

Variables defined this way can be used by any kernel executing on the device (without requiring that they be explicitly passed as a kernel function parameter) and persist for the lifetime of the context, i.e. beyond the lifetime of any single kernel call.

It's not entirely clear from your question, but you can also define an array of data this way.

__device__ float my_variance[32] = {0.0f};

Likewise, allocations created by cudaMalloc live for the duration of the application/context (or until an appropriate cudaFree is encountered) and so there is no need to "pull back the data" to the host if you want to use it in a successive kernel:

float *d_variance;
cudaMalloc((void **)&d_variance), sizeof(float));
my_reduction_kernel<<<...>>>(..., d_variance, ...);
my_thresholding_kernel<<<...>>>(..., d_variance, ...);

Any value set in *d_variance by the reduction kernel above will be properly observed by the thresholding kernel.

Upvotes: 5

Related Questions