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

Question

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

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]; 
    /* ... */ 
} 

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

declaration is incompatible with previous "warp_partial_results" 

сообщение. Что может быть причиной этого?

Answer 1

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

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

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 заголовочном файле они специализируются на некоторых базовых типах, и все.}