Кооперативные группы CUDA: ошибка связывания

Question

После прочтения о кооперативных группах в CUDA 9 я пытался синхронизироваться на уровне сетки.

Я использую Visual Studio 2017, GTX 1060 и CUDA 9.1.

Я изменил свой код следующим образом:

__global__ void ExplicitKernel_American(/* ... */) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    auto grid = cooperative_groups::this_grid();
    if (i < sizeS) {
        //...
        for (int j = 1; j < sizeT; ++j) {
            // ...
            grid.sync(); // __syncthreads();
        }
    }
}

И, как указано в документации, я называю свое ядро ​​следующим образом:

void* Explicit_Args[] = { &PDE_Grid, /* ... */, &sizeS, &sizeT };
cudaLaunchCooperativeKernel(
    (void*)ExplicitKernel_American, 
    dim3((sizeS + TPB - 1) / TPB), 
    dim3(TPB),  
    Explicit_Args
); // TPB being 256...

К сожалению, я получаю ошибки компоновки, как только я добавляю "сетку" в ядро.

error LNK2001: unresolved external symbol __fatbinwrap_38_cuda_device_runtime_compute_70_cpp1_ii_8b1a5d37
fatal error LNK1120: 1 unresolved externals

Я установил -rdc=true и sm_61, но не могу найти, почему он не работает... Есть идеи?

Большое спасибо!

Answer 1

Для использования совместного запуска ядра (кооперативная сетка - CG) требуется графический процессор Pascal или Volta, а также Linux или устройство Windows, работающие в режиме TCC. Если вы проверите deviceProp.cooperativeLaunch свойство в структуре свойств устройства, я думаю, вы обнаружите, что оно не поддерживается на вашем GPU, работающем в режиме WDDM.

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

Однако проблема, о которой вы спрашиваете, - это проблема компиляции / ссылки. Для этого я рекомендую изучить образец кода компьютерной графики (CG), такой как 6_Advanced/reductionMultiBlockCG, В зависимости от того, как вы установите -rdc=true, он может быть неправильно применен к вашему проекту. Правильная методология изложена здесь

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