1

我正在为 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在一个地方更新？谢谢你的帮助。

4

1 回答 1

3

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

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

静态分配是您作为示例给出的，它不起作用：

__shared__ int fitness[16];

在这些情况下，必须在编译时（在设备端代码编译时）知道大小 - 这对您而言并非如此。

使用动态共享内存分配，您无需在内核代码中指定大小 - 您将其留空并添加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博客有一篇不错的文章，其中更深入地介绍了使用共享内存，其中特别介绍了静态和动态共享内存分配。

于 2016-12-30T00:11:48.330 回答