Хранение неиспользованных переменных в CUDA

Question

Я сделал несколько ядер для тестирования пропускной способности, и они не делают никаких полезных вычислений. Минимальный пример -

__global__ void testKernel(float* a) 
{
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    float x;
    x = a[i];
}

Когда я компилирую, я получаю (что неудивительно)

предупреждение: переменная "x" была установлена, но никогда не использовалась

и ядро ​​работает так же быстро, как пустое ядро:

__global__ void donothing() 
{
}

Это указывает на то, что чтение [i] оптимизировано.

Я пробовал такие хитрости, как

volatile float x;

if(x);

(void)(x;)

и они подавляют предупреждение, но ядро ​​все еще слишком быстро завершает работу.

Как я могу убедиться, что бесполезные инструкции действительно выполняются?

Я нашел опцию CU_JIT_OPTIMIZATION_LEVEL, но Google предоставляет в основном ссылки на документацию, а не на то, как ее использовать. Поможет ли мне этот вариант и как мне его использовать?

Answer 1

Попробуйте ввести ветку, в которой хранится переменная:

__global__ void testKernel(float* a, float *b) 
{
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    float x;
    x = a[i];

    if(b)
    {
      *b = x;
    }
}

Стоимость ветки по сравнению со стоимостью передачи памяти незначительна.

На сайте запуска ядра просто передайте нулевой указатель:

testKernel<<<...>>>(a, static_cast<float*>(0));

nvcc не будет выполнять постоянное свертывание при такой степени детализации, поэтому не следует удалять нагрузку, поскольку компилятор не может доказать, что она бесполезна.