【发布时间】：2018-10-04 15:42:28
【问题描述】：
谁能帮助我了解 memCopy2dA 和 memCopy2dB 内核之间的性能差异？
他们应该将大小为 xLen,yLen 的 2D 数据从一个地方复制到另一个地方，但他们使用不同的策略：
当使用 memCopy2dA 时，块/线程覆盖整个 2D 空间，因为该内核假设只复制一个数据点
当使用 memCopy2dB 时，仅为一整 X 行创建块/线程，然后每个内核在 Y 方向上循环以复制所有数据。
根据分析器 (nvvp) 在这两种情况下，GPU 访问内存模式为 100%，并且 X 维度大到足以使“B”内核（Titan X，24SM）的设备饱和。不幸的是，“B”内核速度较慢，在我的机器上结果是：
GB/s: 270.715
GB/s: 224.405
附加问题：是否有可能接近 336.48 GB/s (3505MHz * 384 bits * 2 / 8) 的理论内存带宽限制？至少我的测试显示最大值始终在 271-272 GB/s 左右。
测试代码：
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <chrono>
template<typename T>
__global__ void memCopy2dA(T *in, T *out, size_t xLen, size_t yLen) {
int xi = blockIdx.x * blockDim.x + threadIdx.x;
int yi = blockIdx.y * blockDim.y + threadIdx.y;
if (xi < xLen && yi < yLen) {
out[yi * xLen + xi] = in[yi * xLen + xi];
}
}
template<typename T>
__global__ void memCopy2dB(T *in, T *out, size_t xLen, size_t yLen) {
int xi = blockIdx.x * blockDim.x + threadIdx.x;
if (xi < xLen) {
size_t idx = xi;
for (int y = 0; y < yLen; ++y) {
out[idx] = in[idx];
idx += xLen;
}
}
}
static void waitForCuda() {
cudaDeviceSynchronize();
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(err));
}
int main() {
typedef float T;
size_t xLen = 24 * 32 * 64; //49152
size_t yLen = 1024;
size_t dataSize = xLen * yLen * sizeof(T);
T *dInput;
cudaMalloc(&dInput, dataSize);
T *dOutput;
cudaMalloc(&dOutput, dataSize);
const int numOfRepetitions = 100;
double gigabyte = 1000 * 1000 * 1000;
{
dim3 threadsPerBlock(64, 1);
dim3 numBlocks((xLen + threadsPerBlock.x - 1) / threadsPerBlock.x,
(yLen + threadsPerBlock.y - 1) / threadsPerBlock.y);
auto startTime = std::chrono::high_resolution_clock::now();
for (int i = 0; i < numOfRepetitions; ++i) {
memCopy2dA <<< numBlocks, threadsPerBlock >>> (dInput, dOutput, xLen, yLen);
waitForCuda();
}
auto stopTime = std::chrono::high_resolution_clock::now();
std::chrono::duration<double> elapsed = stopTime - startTime;
std::cout << "GB/s: " << (2 * dataSize * numOfRepetitions) / elapsed.count() / gigabyte << std::endl;
}
{
dim3 threadsPerBlock(64);
dim3 numBlocks((xLen + threadsPerBlock.x - 1) / threadsPerBlock.x);
auto startTime = std::chrono::high_resolution_clock::now();
for (int i = 0; i < numOfRepetitions; ++i) {
memCopy2dB <<< numBlocks, threadsPerBlock >>> (dInput, dOutput, xLen, yLen);
waitForCuda();
}
auto stopTime = std::chrono::high_resolution_clock::now();
std::chrono::duration<double> elapsed = stopTime - startTime;
std::cout << "GB/s: " << ((2 * dataSize * numOfRepetitions) / elapsed.count()) / gigabyte << std::endl;
}
cudaFree(dInput);
cudaFree(dOutput);
return 0;
}
编译：
nvcc -std=c++11 memTest.cu -o memTest
【问题讨论】：
-
第一种方法很可能具有更好的缓存局部性。 gpu 很可能会看到您正在访问顺序内存地址，并且可以使用所需的值预加载内存缓存。由于第二个内核从各处获取内存，因此缓存预加载没有那么大的效果。
-
根据我的理解，两个内核都在做很好的对齐合并内存读/写。在内核“B”中，每个 warp 有 32 个线程处理 4 字节（浮点）元素，一次执行 128 字节的 RD/WR。测试数据维度也“不错” - 一切都非常适合，没有空闲线程左右。这就是我感到困惑的原因。
-
在合并的写入之上可能存在一定程度的内存预取。 GPU 可以假设，由于您刚刚请求了这些字节，您可能很快就会想要它们旁边的字节，因此它可以为您预加载缓存。
-
这是什么 CUDA 版本？关于“附加问题”：
bandwidthTest示例代码通常是这样的复制内核可实现带宽的良好代理。实现峰值理论带宽至少有 2 个限制器。 1. 数据传输有开销——一些位/DRAM 周期通过总线传输，实际上不是用户数据。 2. 像这样的复制内核从读取到写入频繁地“调转总线”。这个周转时间从可实现的带宽中减去。 -
您想验证并尝试使用 gld128 指令来更好地使总线饱和。 32 位负载，即使合并也可能不会使其饱和。此外，使用 restrict 关键字强制（或显式使用）ldg 指令可能有助于使带宽饱和，尤其是在并发写入时。
标签： c++ performance cuda