如何修改CUDA代码以获得100%的GPU负载

How to modify a CUDA code to get 100% GPU load

本文关键字：100% 负载 GPU 代码 何修改 修改 CUDA      更新时间：2023-10-16

如何修改此代码以获得GPU的100%负载？

#include <iostream>
using namespace std;
__global__ void saxpy_parallel(int n, float a, float *x, float *y)
{
// Get the unique ID of this kernel instance
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
{
y[i] = a*x[i] + y[i];
}
}
int main(int argc, char const *argv[])
{
// Tensors length
int const n = 100;
// Define tensors
float x[n], y[n];
for (int i = 0; i < n; ++i)
{
x[i] = 1.0f*i;
y[i] = 1.0f*i;
}
// Device pointers
float *d_x, *d_y;
cudaMalloc(&d_x, n*sizeof(float));
cudaMalloc(&d_y, n*sizeof(float));
if (cudaMemcpy(d_x, &x, n*sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess)
{
printf("Memory Error!n");
return 0;
}
if (cudaMemcpy(d_y, &y, n*sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess)
{
printf("Memory Error!n");
return 0;
}
// Run the kernel
saxpy_parallel<<<4096, 512>>>(n, 2.0, d_x, d_y);
// Retrieve results from the device memory
cudaMemcpy(&y, d_y, n*sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_y);
cudaFree(d_x);
printf("%sn",y[0]);
system("PAUSE");
return 0;
}

好吧，让我们忽略100%GPU负载目标，因为它不现实，也不容易测量。因此，假设您希望优化此代码以更快地运行。目标杠杆是什么？您的算法非常简单，因此不会给自己带来太多机会。然而，我可以看到以下目标

1) 块大小

saxpy_parallel<<<4096, 512>>>

如果512是最好的数字，我会从32或64开始，并在调整内核启动时将大小加倍，以找到该参数的最佳值。

2) 删除不必要的代码

if( i < n )

如果n总是小于i，则可以删除if语句。这可以在内核外部进行控制。可能需要将奇数大小的数组填充为块大小的倍数，以使其发挥作用。

3) 探索矢量类型的使用

CUDA有float2和float4两种类型。因此，重新编写代码以使用这两种方法中的任何一种，并希望通过更少的并行获取和存储以及算术运算来实现更快的内存访问。

4) 解锁环路

每个线程当前正在获取一个x、a和y。尝试获取2、4或8个值

...
y[i] = a*x[i] + y[i];
y[i+1] = a*x[i+1] + y[i+1];
y[i+2] = a*x[i+2] + y[i+2];
y[i+3] = a*x[i+3] + y[i+3];

这需要更少的线程，但每个线程都要做更多的工作。尝试使用2,4,6或8解除干扰价值观

5) 将结果存储到不同的变量中

为结果传入一个额外的参数。然后重新编码

__global__ void saxpy_parallel(int n, float a, float *x, float *y, float * b)
...
b[i] = a*x[i] + y[i];

这就用更多的内存来换取不在同一位置读写。

如果不独立尝试每种方法中的每一种，并在前后测量效果，你将无法确定什么是有效的。那么一些组合可能会更好或者更糟。

试试看，玩得开心，让我们知道！

相关文章：