优化 Cuda 内核时间执行
Optimize Cuda Kernel time execution
我是一名学习Cuda的学生，我想优化我的内核函数的执行时间。结果，我实现了一个计算两张图片之间差异的小程序。所以我比较了 C 中经典 CPU 执行和 Cuda C 中 GPU 执行之间的执行时间。
在这里你可以找到我正在谈论的代码：
int *imgresult_data = (int *) malloc(width*height*sizeof(int));
int size = width*height;
switch(computing_type)
{
case GPU:
HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char)));
HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char)));
HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int)));
HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice));
float time;
cudaEvent_t start, stop;
HANDLE_ERROR( cudaEventCreate(&start) );
HANDLE_ERROR( cudaEventCreate(&stop) );
HANDLE_ERROR( cudaEventRecord(start, 0) );
for(int m = 0; m < nb_loops ; m++)
{
diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res);
}
HANDLE_ERROR( cudaEventRecord(stop, 0) );
HANDLE_ERROR( cudaEventSynchronize(stop) );
HANDLE_ERROR( cudaEventElapsedTime(&time, start, stop) );
HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost));
printf("Time to generate: %4.4f ms \n", time/nb_loops);
break;
case CPU:
clock_t begin = clock(), diff;
for (int z=0; z<nb_loops; z++)
{
// Apply the difference between 2 images
for (int i = 0; i < height; i++)
{
tmp = i*imgresult_pitch;
for (int j = 0; j < width; j++)
{
imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp];
}
}
}
diff = clock() - begin;
float msec = diff*1000/CLOCKS_PER_SEC;
msec = msec/nb_loops;
printf("Time taken %4.4f milliseconds", msec);
break;
}
这是我的内核函数：
__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res)
{
int row = blockIdx.x;
int col = threadIdx.x;
int v = col + row*blockDim.x;
if (row < MAX_H && col < MAX_W)
{
data_res[v] = (int) data2[v] - (int) data1[v];
}
}
我得到了每个执行时间
- CPU：1,3210 毫秒
- GPU：0,3229 毫秒
我想知道为什么 GPU 结果没有达到应有的水平。本人初学Cuda，如有经典错误，还请大家指点。
编辑1：
感谢您的反馈意见。我试图从内核中删除 'if' 条件，但它并没有显着改变我的程序执行时间。
但是，在安装 Cuda 探查器之后，它告诉我我的线程不是 运行 并发的。我不明白为什么我会收到这种消息，但这似乎是真的，因为我使用 GPU 的应用程序只比使用 CPU 快 5 或 6 倍。这个比率应该更大，因为每个线程应该与所有其他线程同时处理一个像素。如果您知道我做错了什么，那将非常有用...
流量。
代码可能还有其他问题，但这是我所看到的。 __global__ void diff 中的以下行被认为不是最佳的：
if (row < MAX_H && col < MAX_W)
{
data_res[v] = (int) data2[v] - (int) data1[v];
}
内核中的条件运算符导致扭曲发散。这意味着 warp 中的 if 和 else 部分是按顺序执行的，而不是并行执行的。此外，您可能已经意识到，if 仅在边界处计算为 false。为避免分歧和不必要的计算，将图像分成两部分：
中心部分，其中 row < MAX_H && col < MAX_W 始终是 true。为此区域创建一个额外的内核。 if在这里是不必要的。
将使用您的 diff 内核的边界区域。
显然，您需要修改调用内核的代码。
在单独的注释中：
GPU 具有面向吞吐量的架构，但不像 CPU 那样面向延迟。这意味着在处理少量数据时，CPU 可能比 CUDA 更快。您是否尝试过使用大型数据集？
CUDA Profiler 是一个非常方便的工具，它会告诉您代码不是最优的。
我认为您没有正确测量时间，内存复制是 GPU 中一个耗时的步骤，您在测量时间时应该考虑到这一点。
我看到了一些你可以测试的细节：
我想你正在使用 MAX_H 和 MAX_H 作为常量，你可以考虑使用 cudaMemcpyToSymbol().
记得使用 __syncthreads() 同步您的线程，这样您就不会在每次循环迭代之间遇到问题。
CUDA 使用扭曲，因此块和每个块的线程数在 8 的倍数时效果更好，但每个块的线程数不得超过 512 个，除非您的硬件支持它。这是每个块使用 128 个线程的示例：<<<(cols*rows+127)/128,128>>>.
还要记住释放您在 GPU 中分配的内存并销毁您创建的时间事件。
在你的内核函数中你可以有一个变量 int v = threadIdx.x + blockIdx.x * blockDim.x .
除了执行时间之外，您是否测试过您的结果是否正确？由于填充，我认为您在处理数组时应该使用 cudaMallocPitch() 和 cudaMemcpy2D()。
您可以做以下两件事来提高 diff 内核的性能：
1。让每个线程做更多的工作
在您的内核中，每个线程只处理一个元素；但是让一个线程做任何事情已经有很多开销，在块和线程级别，包括获取参数、检查条件和进行地址算术。现在，您可以说“哦，但是读取和写入花费的时间要多得多；这个开销可以忽略不计”——但您会忽略这样一个事实，即这些读取和写入的延迟被许多其他的存在所隐藏可以按计划完成工作的 warp。
所以，让每个线程处理多个元素。比方说，4，因为每个线程可以轻松地将 4 个字节一次读入一个寄存器。甚至 8 或 16；用它做实验。当然，您需要相应地调整网格和块参数。
2。 “限制”你的指点
__restrict 不是 C++ 的一部分，但在 CUDA 中受支持。它告诉编译器通过传递给函数的不同指针进行访问永远不会重叠。参见：
- What does the restrict keyword mean in C++?
- Realistic usage of the C99 'restrict' keyword?
使用它允许 CUDA 编译器应用额外的优化，例如loading or storing data via non-coherent cache. Indeed, this happens with your kernel虽然我没有测量过效果。
3。考虑使用“SIMD”指令
CUDA 提供 this 内在：
__device__ unsigned int __vsubss4 ( unsigned int a, unsigned int b )
从 b 中相应的字节值中减去 a 中的每个有符号字节值。如果您可以“接受”结果，而不是期待更大的 int 变量，那可以为您节省一些工作 - 并且随着每个线程元素数量的增加而进展顺利。事实上，它可能会让您进一步增加它以达到最佳效果。
我是一名学习Cuda的学生，我想优化我的内核函数的执行时间。结果，我实现了一个计算两张图片之间差异的小程序。所以我比较了 C 中经典 CPU 执行和 Cuda C 中 GPU 执行之间的执行时间。
在这里你可以找到我正在谈论的代码：
int *imgresult_data = (int *) malloc(width*height*sizeof(int));
int size = width*height;
switch(computing_type)
{
case GPU:
HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char)));
HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char)));
HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int)));
HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice));
HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice));
float time;
cudaEvent_t start, stop;
HANDLE_ERROR( cudaEventCreate(&start) );
HANDLE_ERROR( cudaEventCreate(&stop) );
HANDLE_ERROR( cudaEventRecord(start, 0) );
for(int m = 0; m < nb_loops ; m++)
{
diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res);
}
HANDLE_ERROR( cudaEventRecord(stop, 0) );
HANDLE_ERROR( cudaEventSynchronize(stop) );
HANDLE_ERROR( cudaEventElapsedTime(&time, start, stop) );
HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost));
printf("Time to generate: %4.4f ms \n", time/nb_loops);
break;
case CPU:
clock_t begin = clock(), diff;
for (int z=0; z<nb_loops; z++)
{
// Apply the difference between 2 images
for (int i = 0; i < height; i++)
{
tmp = i*imgresult_pitch;
for (int j = 0; j < width; j++)
{
imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp];
}
}
}
diff = clock() - begin;
float msec = diff*1000/CLOCKS_PER_SEC;
msec = msec/nb_loops;
printf("Time taken %4.4f milliseconds", msec);
break;
}
这是我的内核函数：
__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res)
{
int row = blockIdx.x;
int col = threadIdx.x;
int v = col + row*blockDim.x;
if (row < MAX_H && col < MAX_W)
{
data_res[v] = (int) data2[v] - (int) data1[v];
}
}
我得到了每个执行时间
- CPU：1,3210 毫秒
- GPU：0,3229 毫秒
我想知道为什么 GPU 结果没有达到应有的水平。本人初学Cuda，如有经典错误，还请大家指点。
编辑1： 感谢您的反馈意见。我试图从内核中删除 'if' 条件，但它并没有显着改变我的程序执行时间。
但是，在安装 Cuda 探查器之后，它告诉我我的线程不是 运行 并发的。我不明白为什么我会收到这种消息，但这似乎是真的，因为我使用 GPU 的应用程序只比使用 CPU 快 5 或 6 倍。这个比率应该更大，因为每个线程应该与所有其他线程同时处理一个像素。如果您知道我做错了什么，那将非常有用...
流量。
代码可能还有其他问题，但这是我所看到的。 __global__ void diff 中的以下行被认为不是最佳的：
if (row < MAX_H && col < MAX_W)
{
data_res[v] = (int) data2[v] - (int) data1[v];
}
内核中的条件运算符导致扭曲发散。这意味着 warp 中的 if 和 else 部分是按顺序执行的，而不是并行执行的。此外，您可能已经意识到，if 仅在边界处计算为 false。为避免分歧和不必要的计算，将图像分成两部分：
中心部分，其中
row < MAX_H && col < MAX_W始终是true。为此区域创建一个额外的内核。if在这里是不必要的。将使用您的
diff内核的边界区域。
显然，您需要修改调用内核的代码。
在单独的注释中：
GPU 具有面向吞吐量的架构，但不像 CPU 那样面向延迟。这意味着在处理少量数据时，CPU 可能比 CUDA 更快。您是否尝试过使用大型数据集？
CUDA Profiler 是一个非常方便的工具，它会告诉您代码不是最优的。
我认为您没有正确测量时间，内存复制是 GPU 中一个耗时的步骤，您在测量时间时应该考虑到这一点。
我看到了一些你可以测试的细节：
我想你正在使用 MAX_H 和 MAX_H 作为常量，你可以考虑使用 cudaMemcpyToSymbol().
记得使用 __syncthreads() 同步您的线程，这样您就不会在每次循环迭代之间遇到问题。
CUDA 使用扭曲，因此块和每个块的线程数在 8 的倍数时效果更好，但每个块的线程数不得超过 512 个，除非您的硬件支持它。这是每个块使用 128 个线程的示例：<<<(cols*rows+127)/128,128>>>.
还要记住释放您在 GPU 中分配的内存并销毁您创建的时间事件。
在你的内核函数中你可以有一个变量 int v = threadIdx.x + blockIdx.x * blockDim.x .
除了执行时间之外，您是否测试过您的结果是否正确？由于填充，我认为您在处理数组时应该使用 cudaMallocPitch() 和 cudaMemcpy2D()。
您可以做以下两件事来提高 diff 内核的性能：
1。让每个线程做更多的工作
在您的内核中，每个线程只处理一个元素；但是让一个线程做任何事情已经有很多开销，在块和线程级别，包括获取参数、检查条件和进行地址算术。现在，您可以说“哦，但是读取和写入花费的时间要多得多；这个开销可以忽略不计”——但您会忽略这样一个事实，即这些读取和写入的延迟被许多其他的存在所隐藏可以按计划完成工作的 warp。
所以，让每个线程处理多个元素。比方说，4，因为每个线程可以轻松地将 4 个字节一次读入一个寄存器。甚至 8 或 16；用它做实验。当然，您需要相应地调整网格和块参数。
2。 “限制”你的指点
__restrict 不是 C++ 的一部分，但在 CUDA 中受支持。它告诉编译器通过传递给函数的不同指针进行访问永远不会重叠。参见：
- What does the restrict keyword mean in C++?
- Realistic usage of the C99 'restrict' keyword?
使用它允许 CUDA 编译器应用额外的优化，例如loading or storing data via non-coherent cache. Indeed, this happens with your kernel虽然我没有测量过效果。
3。考虑使用“SIMD”指令
CUDA 提供 this 内在：
__device__ unsigned int __vsubss4 ( unsigned int a, unsigned int b )
从 b 中相应的字节值中减去 a 中的每个有符号字节值。如果您可以“接受”结果，而不是期待更大的 int 变量，那可以为您节省一些工作 - 并且随着每个线程元素数量的增加而进展顺利。事实上，它可能会让您进一步增加它以达到最佳效果。