-1

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

但是，在 CPU 上的所有值之间进行单线程粗暴循环x只会产生比 K520 差 3 倍的性能。我确实知道这是内存密集型任务（线程必须访问和写入他的随机部分summary），但我仍然很难理解 GPU 如何失去它最初的 1000 倍优势。我已经尝试使用内存（因为它在所有线程中都是相同的输入）data以块的形式提供给提要，但没有明显的改进。__constant__nvprof 报告的典型块运行时间为 10-30 秒。

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

编辑：下面是复制问题的示例代码。它可以在 g++（报告运行时间 5s）和 nvcc（报告运行时间 7s）下编译。分析结果如下

==23844== 分析结果：
Time(%) Time Calls Avg Min Max Name
98.86% 4.68899s 1 4.68899s 4.68899s 4.68899s Kernel(Observation*, int*, Info**)
1.09% 51.480ms 4 12.870ms 1.9200us 50.426ms [CUDA memcpy HtoD]
0.06% 2.6634ms 800 3.3290us 3.2950us 5.1200us [CUDA memcpy 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];
}

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

4

1 回答 1

2

您使用的 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 上是昂贵的，为了在内存绑定代码上获得合理的性能，应该避免它们。在这种情况下我们能做些什么呢？在我看来，解决方案是增加更多的并行性。不是每个线程都有一个摘要，而是每个块都有一个摘要。每个线程都可以在 range 的一个子集上工作0...MAX_OBS，并且可以递增应该位于的块范围摘要数组shared memory。在内核结束时，您可以将结果写回全局内存。令人高兴的是，这也解决了您上面提到的缺乏并行性的问题！

接下来是什么？

此时，您应该找到一种方法来衡量还有多少改进空间。你会想知道你离峰值带宽有多近（我发现最好同时考虑你必须移动的数据和你实际移动的数据），如果你仍然很偏离它，你想看看尽可能减少内存访问和进一步优化访问。

于 2014-12-14T17:21:08.943 回答