CUDA中共享内存的重新分配

Question

我有一个关于CUDA C++编程的问题。我正在使用共享内存。但是我需要更大的共享内存。所以我尝试重复使用共享内存。我的代码如下：

__global__ void dist_calculation(...){

   ..........
   {
        //1st pass
       __shared__ short unsigned int shared_nodes[(number_of_nodes-1)*blocksize];

       ............

   }

   {
       //2nd pass
       __shared__ float s_distance_matrix[(number_of_nodes*(number_of_nodes-1))/2];

       ........
   }
}

共享内存无法同时容纳 shared_nodes 和 s_distance_matrix。但是可以分别容纳它们（我已经测试过）。在第二次传递中，程序无法识别 shared_nodes（因为它来自第一次传递），但却显示了一个错误，指出共享内存没有足够的空间。所以看起来，仍然分配了一些空间给 shared_nodes 变量。有没有办法销毁这个分配（比如 cudaFree）？或者有其他建议吗？

- user2026934

---

如果第一遍函数和第二遍函数彼此无关，您可以尝试将这些函数分开。 - ardiyu07

Answer 1

为了执行你的算法，分配一个单一的未类型化缓冲区，足够容纳任何一个数组，并在每次运行算法时重新解释该数组：

__global__ void dist_calculation(...)
{
   const unsigned int num_bytes1 = sizeof(unsigned short) * (number_of_nodes-1) * block_size;

   const unsigned int num_bytes2 = sizeof(float) * (number_of_nodes) * (number_of_nodes-1)) / 2;

   const unsigned int num_shared_bytes = num_bytes1 > num_bytes2? num_bytes1: num_bytes2;

   __shared__ char smem[num_shared_bytes]; 

   unsigned short *shared_nodes = reinterpret_cast<unsigned int*>(smem);
   first_pass(shared_nodes);

   float *distance_matrix = reinterpret_cast<unsigned int*>(smem);
   second_pass(distance_matrix);    
}

- Jared Hoberock