测量 OpenCL 内核的内存吞吐量

Question

我读到了 OpenCL 中的全局内存优化。在其中一张幻灯片中，使用了一个非常简单的内核（如下）来演示内存合并的重要性。

__kernel void measure(__global float* idata, __global float* odata, int   offset) {

    int xid = get_global_id(0) + offset;
    odata[xid] = idata[xid];

}

请看下面我测量内核运行时间的代码

    ret = clFinish(command_queue);
    size_t local_item_size = MAX_THREADS;

    size_t global_item_size = INPUTSIZE;
    struct timeval t0,t1;
    gettimeofday(&t0, 0 );

    //ret = clFinish(command_queue);
    ret = clEnqueueNDRangeKernel(command_queue, measure, 1, NULL,
                                            &global_item_size, &local_item_size, 0, NULL, NULL);




    ret = clFlush(command_queue);
    ret = clFinish(command_queue);

    gettimeofday(&t1,0);
    double elapsed = (t1.tv_sec-t0.tv_sec)*1000000 + (t1.tv_usec-t0.tv_usec);

    printf("time taken = %lf microseconds\n", elapsed);

我传输了大约 0.5 GB 的数据：

#define INPUTSIZE 1024 * 1024 * 128
int main (int argc, char *argv[])
{

   int offset = atoi(argv[1]);
   float* input = (float*) malloc(sizeof(float) * INPUTSIZE); 

现在，结果有点随机。偏移量 = 0，我得到的时间低至 21 微秒。偏移量 = 1 时，我得到的时间范围在 53 到 24400 微秒之间。

谁能告诉我发生了什么事。我认为 offset=0 将是最快的，因为所有线程都将访问连续的位置，因此将发生最少数量的内存事务。

Answer 1

带宽是衡量数据传输速度的指标，在​​这些情况下通常以字节/秒为单位（GPU 内存带宽通常为 GB/秒）。

要计算计算内核的带宽，您只需要知道内核从内存读取/写入内存的数据量，然后将其除以内核执行时间。

您的示例内核让每个工作项（或 CUDA 线程）读取一个浮点数，并写入一个浮点数。如果您启动此内核以复制 2^10 浮点数，那么您将读取 2^10 * sizeof(float) 字节，并写入相同数量（因此总共 8MB）。如果这个内核需要 1ms 来执行，那么你已经达到了8MB / 0.001s = 8GB/s 的带宽。

---

您的新代码 sn-p 显示了您的内核计时方法，表明您只是在计时内核enqueue，而不是运行内核实际花费的时间。这就是为什么您会得到非常低的内核时间 (0.5GB / 0.007ms ~= 71TB/s!)。您应该添加对clFinish() 的调用以获得正确的时间。我通常还会对多次运行进行计时，以使设备预热，这通常会提供更一致的计时：

// Warm-up run (not timed)
clEnqueueNDRangeKernel(command_queue, ...);
clFinish(command_queue);

// start timing
start = ...

for (int i = 0; i < NUM_RUNS; i++)
{
  clEnqueueNDRangeKernel(command_queue, ...);
}
clFinish(command_queue);

// stop timing
end = ...

// Compute time taken, bandwidth etc
average_time = (end-start)/NUM_RUNS;
...

---

评论中的问题：

为什么 offset=0 的性能优于 offset=1,4 或 6？

在 NVIDIA GPU 上，工作项被分组为大小为 32 的“warp”，它们以同步方式执行（其他设备也有类似的方法，只是大小不同）。内存事务与缓存线大小的倍数对齐（例如 64 字节、128 字节等）。考虑一下当 warp 中的每个工作项尝试读取单个 4 字节值时会发生什么（假设它们是连续的，根据您的示例），缓存线大小为 64 字节。

这个 warp 正在读取总共 128 字节的数据。如果这个 128 字节块的开始与 64 字节边界对齐（即如果offset=0），那么这可以在两个 64 字节事务中提供服务。但是，如果此块未与 64 字节边界 (offset=1,4,6,etc) 对齐，则这将需要三个内存事务来获取所有数据。这就是您的性能差异的来源。

如果您将偏移量设置为缓存线大小的倍数（例如 64），那么您可能会获得相当于 offset=0 的性能。