What device is a cuda kernel running on

Question

Is there any way from a kernel (GPU Code) to find out which device the code is running on? You can find blockIdx, threadIdx, blockDim and gridDim, but any way to find the CUDA device id?

For those who wonder, i plan to use it in a OptiX program to have more control over OptiX buffers.

You could certainly discover whatever data you need on the host and then pass it to the kernel when you call it. e.g. `cudaSetDevice(x); mykernel<<<...>>>(x, ...);` — Robert Crovella, Apr 16 '13 at 16:57
I am limited by OptiX API which gives me no control over how the work is split over several GPU's. Still, I would like to determine which device the RT_PROGRAM is running on for memory storage reasons. Therefore, I would need to find this on the device/kernel side, ie. i need a function that's __ device __ tagged. — apartridge, Apr 16 '13 at 16:59
You could use a `__device__` variable to store the device id, since that obviously is per-device... — tera, Apr 16 '13 at 17:02
I cannot since NVIDIA OptiX handles kernel calls and there is no place for me to cudaMemset this variable with the correct device id. But what I can do is to create a OptiX buffer and manually set its content for each device. Its quite a hack but there seems to be no better option. — apartridge, Apr 16 '13 at 17:06

score 1 · Answer 1 · answered Nov 11 '13 at 19:05

This works in OptiX 3.0 and 3.5 but I can't guarantee it will always work. We should provide a real API for this in a future version.

namespace rti_internal_register
{
    __device__ unsigned reg_device_id;
}

rtBuffer<uchar4, 2> output_buffer;

RT_PROGRAM void pinhole_camera()
{
    if( rti_internal_register::reg_device_id == 0 )
        output_buffer[launch_index] = make_color( make_float3( 0.5, 0.0, 0.0 ) );
    else
        output_buffer[launch_index] = make_color( make_float3( 0.0, 0.5, 0.0 ) );
}

What device is a cuda kernel running on

1 Answers1