Использование динамического распределения общей памяти для двух разных векторов

Question

В функции ядра я хочу два вектора разделяемой памяти, как с длиной size (фактически sizeof(float)*size).

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

myKernel<<<numBlocks, numThreads, 2*sizeof(float)*size>>> (...); 

и внутри ядра:

extern __shared__ float row[]; 
extern __shared__ float results[];  

Но это не работает.

Вместо этого я сделал только один вектор extern __shared__ float rowresults[], содержащий все данные, используя выделенную память 2*size. Таким образом, звонки row по-прежнему совпадают, а звонки results похожи на rowresults[size+previousIndex]. И это действительно работает.

Это не большая проблема, потому что я все равно получаю ожидаемые результаты, но есть ли способ разделить мою динамически выделенную разделяемую память на две (или более) разные переменные? Просто для красоты.

Answer 1

C Programming guide section on __shared__ включает примеры, когда вы выделяете несколько массивов с динамически выделяемой совместно используемой памяти:

extern __shared__ float array[]; 
__device__ void func()  // __device__ or __global__ function 
{ 
    short* array0 = (short*)array; 
    float* array1 = (float*)&array0[128]; 
    int* array2 = (int*)&array1[64]; 
} 

Поскольку вы только указатель на элемент и делаете, что новый массив, я полагаю, вы могли бы приспособиться, что использовать динамические смещения вместо статических смещений, которые они имеют в примере. Они также отмечают, что выравнивание должно быть одинаковым, что не должно быть проблемой в вашем случае.