opencl - GPU中缓存未命中的变化

Question

我一直在玩一个OpenCL内核，该内核可以访问7个全局内存缓冲区，对值进行某些处理，然后将结果存储回第8个全局内存缓冲区。正如我观察到的，随着输入大小的增加，L1缓存未命中率（=丢失（丢失+命中））变化很大。我找不到这种变化的来源。这里的输入大小表示全局工作项的数量（2的幂，是工作组大小的倍数）。工作组大小的数量保持为256。

这些就是结果。这些显示L1高速缓存未命中率。从4096个工作项（16个工作组）开始。

0.677125
0.55946875
0.345994792
0.054078125
0.436167969
0.431871745
0.938546224
0.959258789
0.952941406
0.955016479

探查器说它每个线程使用18个寄存器。
这是代码（函数TTsum（）应该只执行一堆相关的超越操作，所以我认为它与缓存无关）：

float TTsum(float x1, float x2, float x3, float x4, float x5, float x6, float x7)
{
        float temp = 0;
        for (int j = 0; j < 2; j++)
                temp = temp +  x1 + (float)x2 + x3 + x4 + x5 + x6 + x7;
        temp = sqrt(temp);
        temp = exp(temp);
        temp = temp / x1;
        temp = temp / (float)x2;
        for (int j = 0; j < 20; j++) temp = sqrt(temp);
        return temp;
}

__kernel void histogram(__global float* x1,
                        __global int* x2,
                        __global float* x3,
                        __global float* x4,
                        __global float* x5,
                        __global float* x6,
                        __global float* x7,
                        __global float* y)
{
  int id = get_global_id(0);    
  for (int j = 0; j < 1000; j++)
    y[id] = TTsum(x1[id], x2[id], x3[id], x4[id], x5[id], x6[id], x7[id]);
}

有人可以解释缓存行为吗？实验在GTX580中完成。

Answer 1

在CUDA中计算直方图非常困难。我相信y []的随机访问很可能是您观察到的行为的原因。如果您还没有的话，请阅读以下内容：http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/histogram256/doc/histogram.pdf