Получение ошибки CUDA "объявление несовместимо с предыдущим"variable_name"

Question

Я пытаюсь скомпилировать программу, включая ядро ​​с MSVS 2012 и CUDA. Я использую разделяемую память, но в отличие от этого вопроса, касающегося той же проблемы, я использую имя моей переменной для разделяемой памяти этого ядра только один раз, поэтому переопределения не возникает. С таким кодом:

template<typename T>
__global__ void mykernel(
    const T* __restrict__ data,
    T*       __restrict__ results) 
{
    extern __shared__ T warp_partial_results[];
    /* ... */
    warp_partial_results[lane_id] = something;
    /* ... */
    results[something_else] = warp_partial_results[something_else];
    /* ... */
}

который создается для нескольких типов (например, float, int, unsigned int), я получаю страшно

declaration is incompatible with previous "warp_partial_results"

сообщение. Что может вызвать это?

Answer 1

CUDA не сразу "поддерживает" динамически распределяемые массивы совместно используемой памяти в шаблонных функциях, поскольку (по-видимому) генерирует фактические определения этих extern'ов. Если вы создадите шаблонную функцию для нескольких типов, определения будут конфликтовать.

Обходной путь доступен в форме специализации шаблона через классы. Увидеть:

http://www.ecse.rpi.edu/~wrf/wiki/ParallelComputingSpring2015/cuda/nvidia/samples/0_Simple/simpleTemplates/sharedmem.cuh

Вы используете обходной путь следующим образом:

template<class T> __global__ void foo( T* g_idata, T* g_odata)
{
    // shared memory
    // the size is determined by the host application

    SharedMem<T> shared;
    T* sdata = shared.getPointer();

    // .. the rest of the code remains unchanged!
}

getPointer() имеет ^* специализированную реализацию для каждого типа, которая возвращает другой указатель, например extern __shared__ float* shared_mem_float или же extern __shared__ int* shared_mem_int и т.п.

_{(*) - На самом деле, нет. в поставляемом заголовочном файле NVidia они специализируются на некоторых основных типах, и это все.}