Освобождение от общей разделяемой памяти CUDA

Question

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

__shared__ unsigned char flag; 
__shared__ unsigned int values [ BLOCK_DIM ]; 

Позже в коде, мне нужно определить массив, с рассмотрением ранее определенной совместно используемой памяти превышает общий предел памяти для установки блока. Как я могу это сделать без грязных работ по повторному использованию ранее определенной разделяемой памяти? Или NVCC достаточно умен, чтобы распознавать зависимости по трассировке ядра и освобождает его когда-либо, используя общие переменные? Мое устройство - GeForce GTX 780 (CC = 3,5).

Answer 1

В C/C++ невозможно освободить статически определенные массивы.

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

extern __shared__ float foo[]; 

в функции ядра и запустить вашу функцию ядра как

myKernel<<<numBlocks, numThreads, sh_mem_size>>> (...); 

Помните, что вы можете управлять несколькими массивами, играя с указателями. Посмотрите на CUDA C Programming Guide для получения дополнительной информации. Например, со ссылкой на руководство

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]; 
} 

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