测量 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 usecs 到 24400 usecs 之间。

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

Answer 1

带宽是衡量数据传输速度的指标，在这些情况下通常以 bytes/second 衡量（GPU 内存带宽通常 GB/s）。

要计算计算内核的带宽，您只需要知道内核 reads/writes from/to 内存有多少数据，然后将其除以内核执行所需的时间。

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

的带宽

显示您的内核计时方法的新代码片段表明您只是在计时内核入队，而不是运行核心。这就是为什么您获得非常低的内核时序 (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;
...

来自评论的问题：

Why does offset=0 perform better than offset=1,4 or 6?

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

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

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

的性能

测量 OpenCL 内核的内存吞吐量

Measuring OpenCL kernel's memory throughput

profiling

opencl