测量 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 微秒。offset = 1 时，我得到的时间范围在 53 到 24400 微秒之间。

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

Answer 1

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

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

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

显示内核计时方法的新代码片段表明您只计时内核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 的“扭曲”中，它们以锁步方式执行（其他设备也有类似的方法，只是大小不同）。内存事务与缓存行大小的倍数对齐（例如 64 字节、128 字节等）。考虑当 warp 中的每个工作项尝试读取单个 4 字节值（假设它们是连续的，根据您的示例）时会发生什么，缓存行大小为 64 字节。

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

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