c++ - 我如何在我的 CUDA 内核中使用共享内存？

Question

我有以下 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 我想找到一种在共享 (__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);
}

Answer 1

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

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

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

但是您的每个数据都很大(838 维 vector )，也许策略也很大。您可能无法在共享内存中缓存很多(每个 block 只有 ~48k，具体取决于设备类型)。所以情况变成了矩阵-矩阵乘法之类的东西。为此，您可能会从以下链接中的矩阵乘法代码中获得一些提示。

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