如何在CUDA内核中使用共享内存

How can I use shared memory here in my CUDA kernel?

本文关键字：共享 内存 内核 CUDA      更新时间：2023-10-16

我有以下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__)内存中的方法。我不明白的是如何使用多个线程来填充共享变量。

我见过这样的例子，其中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]
...

无限制，数据是一个固定大小的2D阵列，类似于：

[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);
}

对于未来寻求类似答案的人来说，以下是我为内核函数所做的：

__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。我不知道事情是否应该这样做，但这对我来说是有效的。

正如您的评论中所确认的，您希望对每一个3.5米的数据应用20k(这个数字来自您上一个问题)策略，并检查20kx3.5米的结果。

如果没有共享内存，您必须从全局内存中读取2万次所有数据或350万次所有策略。

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

然而，你的每一个数据都很大(838-D矢量)，也许策略也很大。您可能无法在共享内存中缓存大量内存(根据设备类型，每个块只能缓存~48k)。所以情况就变成了矩阵乘法。为此，您可能会从以下链接中的矩阵乘法代码中得到一些提示。

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

相关文章：