在这种情况下，有什么技巧可以提高CUDA性能，例如声明全局/局部变量，参数传递，内存复制。
我试图在下面的示例中找出sum_gpu_FAST和sum_gpu_SLOW之间的两个性能差异太大的原因。
在这里，您可以看到整个示例代码。
#include <iostream>
#include <chrono>
#define N 10000000
__global__
void sum_gpu_FAST(int (&data)[N][2], int& sum, int n) { // runtime : 2.42342s
int s = 0;
for (int i = 0; i < n; i++)
s += data[i][0] * 10 + data[i][1];
sum = s;
}
__global__
void sum_gpu_SLOW(int (&data)[N][2], int& sum, int n) { // runtime : 436.64ms
sum = 0;
for (int i = 0; i < n; i++) {
sum += data[i][0] * 10 + data[i][1];
}
}
void sum_cpu(int (*data)[2], int& sum, int n) {
for (int i = 0; i < n; i++) {
sum += data[i][0] * 10 + data[i][1];
}
}
int main()
{
int (*v)[2] = new int[N][2];
for (int i = 0; i < N; i++)
v[i][0] = 1, v[i][1] = 3;
printf ("-CPU------------------------------------------------\n");
{
int sum = 0;
auto start = std::chrono::system_clock::now();
sum_cpu(v, sum, N);
auto end = std::chrono::system_clock::now();
// print output
std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
}
printf ("-GPU-Ready------------------------------------------\n");
int *dev_sum = nullptr;
int (*dev_v)[N][2] = nullptr;
cudaMalloc((void **)&dev_v, sizeof(int[N][2]));
cudaMalloc((void **)&dev_sum, sizeof(int));
cudaMemcpy(dev_v, v, sizeof(int[N][2]), cudaMemcpyHostToDevice);
printf("-GPU-FAST-------------------------------------------\n");
{
int sum = 0;
auto start = std::chrono::system_clock::now();
sum_gpu_FAST<<<1, 1>>> (*dev_v, *dev_sum, N);
cudaDeviceSynchronize(); // wait until end of kernel
auto end = std::chrono::system_clock::now();
// print output
cudaMemcpy( &sum, dev_sum, sizeof(int), cudaMemcpyDeviceToHost );
std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
}
printf("-GPU-SLOW-------------------------------------------\n");
{
int sum = 0;
auto start = std::chrono::system_clock::now();
sum_gpu_SLOW<<<1, 1>>> (*dev_v, *dev_sum, N);
cudaDeviceSynchronize(); // wait until end of kernel
auto end = std::chrono::system_clock::now();
// print output
cudaMemcpy( &sum, dev_sum, sizeof(int), cudaMemcpyDeviceToHost );
std::cout << sum << " / " << (end-start).count() / 1000000 << "ms" << std::endl;
}
printf("----------------------------------------------------\n");
return 0;
}
最佳答案
我试图在下面的示例中找出sum_gpu_FAST和sum_gpu_SLOW之间的两个性能差异太大的原因。
在快速情况下，您正在创建一个（大概）包含在寄存器中的局部变量：
int s = 0;
在循环迭代期间，从全局存储器进行读取，但是唯一的写操作是对寄存器的写入：
for (int i = 0; i < n; i++)
s += data[i][0] * 10 + data[i][1];
在慢速情况下，运行总和包含在全局内存中的一个变量中：
sum = 0;
因此，在每次循环迭代时，更新后的值都会写入全局内存：
for (int i = 0; i < n; i++) {
sum += data[i][0] * 10 + data[i][1];
因此，循环在每次迭代时都有额外的开销要写入全局存储器，这比将和保持在寄存器中要慢。
我不会完全剖析SASS代码来比较这两种情况，因为编译器会在快速情况下围绕循环展开和可能的其他因素做出其他决策，但是我的猜测是不需要将结果存储到循环迭代过程中的全局内存也大大有助于循环展开。但是，对于每种情况，我们都可以根据SASS代码的结尾进行简单的推论：
Function : _Z12sum_gpu_FASTRA10000000_A2_iRii
.headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
/*0000*/ MOV R1, c[0x0][0x28] ; /* 0x00000a0000017a02 */
/* 0x000fd00000000f00 */
...
/*0b00*/ STG.E.SYS [R2], R20 ; /* 0x0000001402007386 */
/* 0x000fe2000010e900 */
/*0b10*/ EXIT ; /* 0x000000000000794d */
/* 0x000fea0003800000 */
在上面的快速案例中，我们看到在内核的末尾，在return语句（
STG）之前，并且在内核中的任何循环之外，只有一条全局存储（EXIT）指令。尽管我还没有显示全部，但快速内核中确实没有其他STG指令，除了最后一个。对于慢速内核的末尾，我们看到了另一个故事： code for sm_70
Function : _Z12sum_gpu_SLOWRA10000000_A2_iRii
.headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
/*0000*/ IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ; /* 0x00000a00ff017624 */
/* 0x000fd000078e00ff */
...
/*0460*/ STG.E.SYS [R2], R7 ; /* 0x0000000702007386 */
/* 0x0005e2000010e900 */
/*0470*/ @!P0 BRA 0x2f0 ; /* 0xfffffe7000008947 */
/* 0x000fea000383ffff */
/*0480*/ EXIT ; /* 0x000000000000794d */
/* 0x000fea0003800000 */
慢速内核在循环内部以
STG指令结束循环。慢速内核在整个内核中还具有许多STG指令的实例，这可能是由于编译器展开所致。关于c++ - CUDA的性能取决于声明变量，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/59385108/