如何在我的 CUDA 内核中使用共享内存？答案

【问题标题】：How can I use shared memory here in my CUDA kernel?如何在我的 CUDA 内核中使用共享内存？
【发布时间】：2016-06-07 17:08:44
【问题描述】：

我有以下 CUDA 内核：

__global__ void optimizer_backtest(double *data, Strategy *strategies, int strategyCount, double investment, double profitability) {
    // Use a grid-stride loop.
    // Reference: https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < strategyCount;
         i += blockDim.x * gridDim.x)
    {
        strategies[i].backtest(data, investment, profitability);
    }
}

TL;DR 我想找到一种方法将data 存储在共享 (__shared__) 内存中。我不明白的是如何使用多个线程填充共享变量。

我见过像this one 这样的示例，其中data 被逐个线程复制到共享内存中（例如myblock[tid] = data[tid]），但我不确定在我的情况下如何执行此操作。问题是每个线程都需要访问整个“行”（展平）数据，每次迭代通过数据集（请参阅下文中调用内核的位置）。

我希望是这样的：

__global__ void optimizer_backtest(double *data, Strategy *strategies, int strategyCount, int propertyCount, double investment, double profitability) {
    __shared__ double sharedData[propertyCount];

    // Use a grid-stride loop.
    // Reference: https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < strategyCount;
         i += blockDim.x * gridDim.x)
    {
        strategies[i].backtest(sharedData, investment, profitability);
    }
}

这里有更多详细信息（如果需要更多信息，请询问！）：

strategies 是指向Strategy 对象列表的指针，data 是指向分配的扁平数据数组的指针。

在backtest() 中，我像这样访问数据：

data[0]
data[1]
data[2]
...

未展平，数据是一个固定大小的二维数组，类似这样：

[87.6, 85.4, 88.2, 86.1]
 84.1, 86.5, 86.7, 85.9
 86.7, 86.5, 86.2, 86.1
 ...]

至于内核调用，我对数据项进行了迭代，对n个数据行（约350万）调用了n次：

int dataCount = 3500000;
int propertyCount = 4;

for (i=0; i<dataCount; i++) {
    unsigned int dataPointerOffset = i * propertyCount;

    // Notice pointer arithmetic.
    optimizer_backtest<<<32, 1024>>>(devData + dataPointerOffset, devStrategies, strategyCount, investment, profitability);
}

【问题讨论】：

为什么sharedData 是一个指针数组？完全不清楚您在这里要做什么
它不是一个指针数组。这相当于指向内存中数组的指针。正如我已经解释过的，它是一个像这样的扁平结构：[87.6, 85.4, 88.2, 86.1, 84.1, 86.5, 86.7, 85.9, 86.7, 86.5, 86.2, 86.1, ...]。 data[1] 是 85.4。这正是Strategy::backtest() 使用data 的方式：它通过索引访问项目。 data初始化如下：double *data = (double*)malloc(dataPointCount * dataPropertyCount * sizeof(double));我不知道还有什么要解释的。
哦，我明白你在说什么。 sharedData 应该是 __shared__ double sharedData[propertyCount]。我会更新这个。谢谢你。请不要因为一个错字而投票结束我的问题。
您真的是说您不能编写代码来将 4 个双精度值加载到共享内存数组中吗？究竟你这样做的问题是什么？我不会因为拼写错误而投票关闭它，我投票关闭它是因为它是一个模糊、广泛的帖子，有很多词，没有真正的具体问题。
有点不清楚你想用你的代码做什么，为什么你需要共享内存。如果你能稍微解释一下这些会有所帮助。共享内存通常用于跨线程交换数据。所以我猜你想对 3.5m 数据中的每一个应用 20k 策略并检查 20k x 3.5m 结果？如果是，那么使用共享内存听起来是个不错的情况。但是您需要比 4 个元素更大的共享内存，并且您可能还希望使用共享内存缓存策略。

标签： c++ cuda shared-memory

【解决方案1】：

正如您在评论中确认的那样，您希望对 350 万个数据中的每一个应用 20k（这个数字来自您之前的问题）策略并检查 20k x 3.5m 的结果。

如果没有共享内存，您必须从全局内存中读取所有数据 20k 次或所有策略 3.5m 次。

共享内存可以通过减少全局内存访问来加速您的程序。假设您每次可以读取 1k 个策略和 1k 个数据到共享内存，检查 1k x 1k 个结果，然后重复此操作，直到全部检查完。通过这种方式，您可以将全局 mem 访问减少 20 次所有数据和 3.5k 次所有策略。这种情况类似于向量-向量叉积。您可以找到一些参考代码以获取更多详细信息。

但是，您的每个数据都很大（838-D 向量），也许策略也很大。您可能无法在共享内存中缓存很多它们（每个块仅约 48k，具体取决于设备类型）。所以情况变成了矩阵-矩阵乘法之类的东西。为此，您可能会从以下链接中的矩阵乘法代码中获得一些提示。

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory

【讨论】：

【解决方案2】：

对于未来寻找类似答案的人，这是我最终得到的内核函数：

__global__ void optimizer_backtest(double *data, Strategy *strategies, int strategyCount, double investment, double profitability) {
    __shared__ double sharedData[838];

    if (threadIdx.x < 838) {
        sharedData[threadIdx.x] = data[threadIdx.x];
    }

    __syncthreads();

    // Use a grid-stride loop.
    // Reference: https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < strategyCount;
         i += blockDim.x * gridDim.x)
    {
        strategies[i].backtest(sharedData, investment, profitability);
    }
}

请注意，我在我的应用程序中同时使用了 .cuh 和 .cu 文件，并将其放在 .cu 文件中。另请注意，在编译目标文件时，我在 Makefile 中使用了--device-c。我不知道事情是否应该这样做，但这对我有用。

【讨论】：