我有一个数据处理任务,可以通过以下方式进行风格化。我有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
中的块范围的摘要数组。在内核结束时,您可以将结果写回全局内存。令人高兴的是,这也解决了上面提到的缺乏并行性的问题!
接下来呢?
在这一点上,你应该找到一种方法来衡量有多少改进空间。您需要计算出您离峰值带宽有多近(我发现最好同时考虑您必须移动的数据和您实际移动的数据(,如果您仍然明显偏离峰值带宽,您希望查看减少内存访问和进一步优化访问,如果可能的话。