CUDA 在哪里声明共享内存分配的常量

Question

我正在为 1024 个矩阵运行适应度函数，每个矩阵都有自己的块并且大小相同。每个块都有n*n 线程（矩阵的维度）并且需要有n*n 共享内存，这样我就可以轻松地进行求和。但是，所有矩阵的维度n 在运行前是可变的（即可以手动更改，尽管始终是 2 的幂，所以求和很简单）。这里的问题是必须使用常量分配共享内存，但我还需要将值从主机传递给内核。我在哪里声明维度n，以便它对 CPU 可见（用于传递给内核）并且可用于声明共享内存的大小（在内核内）？

我的代码结构如下：

来自main.cu我调用内核：

const int num_states = 1024
const int dimension = 4

fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return);

然后在kernel.cu 我有：

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return) {
    __shared__ int fitness[16]; <<-- needs to be dimension * dimension
    //code
}

numbers 是一个表示 1024 个矩阵的数组，dimension 是行列长度，num_states 是 1024，fitness_return 是一个长度为 1024 的数组，用于保存每个矩阵的适应度值。在内核中，共享内存是用dimension 的平方硬编码的（所以在这个例子中dimension 是4）。

我在哪里以及如何声明dimension，以便它可以用于分配共享内存以及调用内核，这样我只需在一个地方更新dimension？感谢您的帮助。

Answer 1

分配的共享内存量在所有块上是统一的。您可能在每个块中使用不同数量的共享内存，但它仍然可用。此外，无论如何共享内存的数量都相当有限，因此 n*n 元素不能超过maximum amount of space（通常为 48KiB）；对于float-type 元素（每个 4 字节），这意味着 n

现在，有两种分配共享内存的方法：静态和动态。

静态分配是您作为示例给出的，这是行不通的：

__shared__ int fitness[16];

在这些情况下，大小必须在编译时知道（在设备端代码编译时）——这对你来说不是这样。

使用动态共享内存分配，您不在内核代码中指定大小 - 您在leave it empty 和前面加上extern：

extern __shared__ int fitness[];

相反，你在启动内核时指定数量，不同块的线程不一定知道它是什么。

但在您的情况下，线程 确实 需要知道 n 是什么。好吧，只需将其作为内核参数传递即可。所以，

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return,
    unsigned short fitness_matrix_order /* that's your n*/) 
{
    extern __shared__ int fitness[];
    /* ... etc ... */
}

nVIDIA 的Parallel-for-all 博客有a nice post 更深入地介绍了使用共享内存，具体包括静态和动态共享内存分配。