我了解了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将是最快的，因为所有线程都将访问连续的位置，因此将发生最少数量的内存事务。
带宽是数据传输速度的一种度量，在这些情况下通常以字节/秒为单位进行度量（GPU内存带宽通常为GB / s）。
要计算计算内核的带宽，您只需要知道内核从内存读取/向内存写入的数据量，然后将其除以内核执行所需的时间即可。
您的示例内核使每个工作项（或CUDA线程）读取一个浮点数，并写入一个浮点数。如果启动此内核以复制2^10浮点数，那么您将读取2^10 * sizeof(float)字节并写入相同的数量（8MB总计）。如果该内核执行需要1毫秒的时间，则您已达到的带宽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;
...
来自评论的问题：
为什么offset = 0的表现要好于offset = 1,4或6？
在NVIDIA GPU上，工作项目分为大小为32的“变形”，这些变形以锁步的方式执行（其他设备具有类似的方法，只是大小不同）。内存事务与高速缓存行大小的倍数对齐（例如64字节，128字节等）。考虑一下当扭曲中的每个工作项都尝试读取单个4字节值（假设您的示例是连续的）时，缓存行大小为64字节时会发生什么。
此扭曲总共读取128个字节的数据。如果此128字节块的开头与64字节边界对齐（即if offset=0），则可以在两个64字节事务中进行处理。但是，如果该块未与64字节边界（offset=1,4,6,etc）对齐，则将需要三个内存事务来获取所有数据。这就是您的性能差异的来源。
如果将偏移量设置为缓存行大小的倍数（例如64），则可能会获得与相等的性能offset=0。
本文收集自互联网，转载请注明来源。
如有侵权，请联系 [email protected] 删除。
我来说两句