优化 CUDA 的特定内存使用情况



我有一个数据处理任务,可以通过以下方式进行风格化。我有data(~1-10GB(和一个函数,该函数基于此data和一些(双倍(输入x生成summary(~1MB(。我需要为 ~1000 个x值获取此summary,这对于 GPU 来说似乎是一项完美的任务。重复一遍,所有线程的输入data都相同,并且以线性方式读取,但每个线程必须产生自己的summary。函数针对不同的x独立执行。

但是,在 CPU 上暴力单线程循环x的所有值只会产生比 K520 差 3 倍的性能。我确实知道这是内存密集型任务(线程必须访问并写入他summary的随机部分(,但我仍然难以理解 GPU 如何失去其最初的 1000 倍优势。我尝试使用内存将data馈送到块__constant__(因为它在所有线程中都是相同的输入(,但没有明显的改进。根据 nvprof 的报告,典型的块运行时间为 10-30 秒。

我将不胜感激对适合此任务的优化的任何见解。

编辑:下面是复制问题的示例代码。它可以在 g++(报告运行时间为 5 秒(和 nvcc(报告运行时间为 7 秒(下编译。性能分析结果如下

==23844== 性能分析结果:
时间(%( 通话时间 平均最小最大名称
98.86% 4.68899s 1 4.68899s 4.68899s 4.68899s 内核(观察*, int*, 信息**(
1.09% 51.480毫秒 4 12.870毫秒 1.9200us 50.426毫秒 [CUDA 内存 HtoD]
0.06% 2.6634毫秒 800 3.3290美国 3.2950美国 5.1200us [CUDA 内存 DtoD]
0.00% 4.3200us 1 4.3200us 4.3200us 4.3200us [CUDA memcpy DtoH]

#include <iostream>
#include <fstream>
#include <cstdlib>
#include <ctime>
#include <cstring>
#define MAX_OBS 1000000
#define MAX_BUCKETS 1000
using namespace std;
// Cross-arch defines
#ifndef __CUDACC__
#define GPU_FUNCTION
#define cudaSuccess 0
typedef int cudaError_t;
struct dim3
{
    int x;
    int y;
    int z;
} blockIdx, threadIdx;
enum cudaMemcpyKind
{
    cudaMemcpyHostToDevice = 0,
    cudaMemcpyDeviceToHost = 1, 
    cudaMemcpyDeviceToDevice = 2
};
cudaError_t cudaMalloc(void ** Dst, size_t bytes)
{
    return !(*Dst = malloc(bytes));
}
cudaError_t cudaMemcpy(void * Dst, const void * Src, size_t bytes, cudaMemcpyKind kind)
{
    return !memcpy(Dst, Src, bytes);
}
#else
#define GPU_FUNCTION __global__
#endif
// Basic observation structure as stored on disk
struct Observation
{
    double core[20];
};
struct Info
{
    int left;
    int right;
};
GPU_FUNCTION void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[threadIdx.x * 40 + threadIdx.y];
    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_obs[i].core[threadIdx.x] < (threadIdx.x + 1) * threadIdx.y)
            summary[d_bucket[i]].left++;
        else
            summary[d_bucket[i]].right++;
    }
}
int main()
{
    srand((unsigned int)time(NULL));
    // Generate dummy observations
    Observation * obs = new Observation [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        for (int j = 0; j < 20; j++)
            obs[i].core[j] = (double)rand() / RAND_MAX;
    // Attribute observations to one of the buckets
    int * bucket = new int [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        bucket[i] = rand() % MAX_BUCKETS;
    Info summary[MAX_BUCKETS];
    for (int i = 0; i < MAX_BUCKETS; i++)
        summary[i].left = summary[i].right = 0;
    time_t start;
    time(&start);
    // Init device objects
    Observation * d_obs;                    
    int * d_bucket; 
    Info * d_summary;
    Info ** d_summaries;
    cudaMalloc((void**)&d_obs, MAX_OBS * sizeof(Observation));
    cudaMemcpy(d_obs, obs, MAX_OBS * sizeof(Observation), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_bucket, MAX_OBS * sizeof(int));
    cudaMemcpy(d_bucket, bucket, MAX_OBS * sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_summary, MAX_BUCKETS * sizeof(Info));
    cudaMemcpy(d_summary, summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyHostToDevice);
    Info ** tmp_summaries = new Info * [20 * 40];
    for (int k = 0; k < 20 * 40; k++)           
        cudaMalloc((void**)&tmp_summaries[k], MAX_BUCKETS * sizeof(Info));
    cudaMalloc((void**)&d_summaries, 20 * 40 * sizeof(Info*));
    cudaMemcpy(d_summaries, tmp_summaries, 20 * 40 * sizeof(Info*), cudaMemcpyHostToDevice);
    for (int k = 0; k < 20 * 40; k++)
        cudaMemcpy(tmp_summaries[k], d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToDevice);
#ifdef __CUDACC__
    Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);
#else
    for (int k = 0; k < 20 * 40; k++)
    {
        threadIdx.x = k / 40;
        threadIdx.y = k % 40;
        Kernel(d_obs, d_bucket, d_summaries);
    }
#endif      
    cudaMemcpy(summary, d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToHost);
    time_t end;
    time(&end);
    cout << "Finished calculations in " << difftime(end, start) << "s" << endl;
    cin.get();
    return 0;
}

编辑2:我尝试通过并行处理艰难的分散内存访问来重新设计代码。简而言之,我的新内核看起来像这样

__global__ void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         double * values,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[blockIdx.x * 40 + blockIdx.y];
    __shared__ Info working_summary[1024];
    working_summary[threadIdx.x] = summary[threadIdx.x];
    __syncthreads();
    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_bucket[i] != threadIdx.x) continue;
        if (d_obs[i].core[blockIdx.x] < values[blockIdx.y])
            working_summary[threadIdx.x].left++;
        else
            working_summary[threadIdx.x].right++;
    }
    __syncthreads();
    summary[threadIdx.x] = working_summary[threadIdx.x];
} 

<<<dim(20, 40, 1), 1000>>>需要 18 秒,<<<dim(20,40,10), 1000>>> ---需要 172 秒,这比单个 CPU 线程更差,并且并行任务的数量呈线性增加。

您使用的 K520 板有两个 GPU,每个 GPU 都有 8 个流式多处理器,我相信每个 GPU 的峰值带宽为 ~160 GB/s。使用上面的代码,您应该受到此带宽的限制,并且应该考虑每个GPU至少获得100 GB/s(尽管我的目标是单个GPU开始(。也许你无法击中它,也许你会击败它,但这是一个很好的目标。

块数

首先要做的是修复您的启动参数。这一行:

Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);

意味着您正在启动 1 个包含 800 个线程的 CUDA 块。这对于 GPU 来说远远不够并行。您至少需要与流式多处理器一样多的块(即 8 个(,最好是更多(即 100+(。这将给您带来很大的性能改进。800 路并行性对于 GPU 来说是不够的。

分散写入

GPU 对访问模式可能相当敏感。以下代码:

summary[d_bucket[i]].left++;

是否将分散的 4 字节写入摘要。分散的内存事务在 GPU 上是昂贵的,为了在内存绑定代码上获得合理的性能,应避免使用它们。在这种情况下,我们能做些什么呢?在我看来,解决方案是增加更多的并行性。与其让每个线程有一个摘要,不如为每个有一个摘要。每个线程都可以处理范围0...MAX_OBS的子集,并且可以递增应该位于shared memory中的块范围的摘要数组。在内核结束时,您可以将结果写回全局内存。令人高兴的是,这也解决了上面提到的缺乏并行性的问题!

接下来呢?

在这一点上,你应该找到一种方法来衡量有多少改进空间。您需要计算出您离峰值带宽有多近(我发现最好同时考虑您必须移动的数据和您实际移动的数据(,如果您仍然明显偏离峰值带宽,您希望查看减少内存访问和进一步优化访问,如果可能的话。

相关内容

  • 没有找到相关文章

最新更新