如何修改 CUDA 代码以获得 100% GPU 负载 [关闭]答案

【问题标题】：How to modify a CUDA code to get 100% GPU load [closed]如何修改 CUDA 代码以获得 100% GPU 负载 [关闭]
【发布时间】：2014-01-26 15:53:57
【问题描述】：

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

#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("%s\n",y[0]);

    system("PAUSE");
    return 0;
}

【问题讨论】：

定义“100% GPU 使用率”——什么意思？
@talonmies 对不起。我的意思是：legitreviews.com/images/reviews/1688/GPUzLoad.png看看声音“GPU Load”
您说的传感器参数GPU负载恐怕与CUDA编程无关。
我的意思正是@talonmies 在他的评论中的意思。我在 CUDA C Programming Guide 和 CUDA C Best Practices Guide 中都找不到 GPU load 的定义。所以我认为你必须向这个社区解释一些事情。如果我们不知道它的定义，如何最大化您的 GPU 负载 参数？投票结束。
我还没有说你链接的问题很清楚。 GPU 使用的定义对您来说似乎不是很清楚，因为您还不能提供它。说起来可能不愉快，但这应该会让你认为你正在处理的问题还没有以令人满意的方式正式制定。

标签： c++ cuda gpgpu gpu

【解决方案1】：

好的，让我们忽略 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];

这会以更多的内存换取不读取和写入同一位置。

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

试试看，玩得开心，告诉我们！

【讨论】：