GPU缓存未命中的变化

18

我一直在实验一个访问7个全局内存缓冲区的OpenCL内核,在值上执行某些操作,并将结果存储回第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上进行的。


你应该向TTsum展示代码,因为如果它的变量没有保存在寄存器中,它也将使用L2缓存。 - talonmies
我修改了问题。你是说L2用于寄存器溢出吗?我以为是L1?顺便提一下,这里的所有值都只针对L1缓存。 - Zk1001
我猜你正在使用随机输入数据。我建议尝试不使用exp和sqrt调用的内核。Exp是一个函数,可能会使用线程本地内存,并且根据输入而定,其运行时间不会恒定。如果内核被限制为一系列常量时间、基于寄存器的乘加操作序列,则可能会看到完全不同的结果。 - talonmies
好的,我将尝试仅使用MADDs。您怀疑那些超越函数使用全局内存吗?这怎么可能呢? - Zk1001
我在测试中删除了所有的exp和sqrt,只保留了add(第一个for循环)。但缓存行为是相似的。不仅缓存未命中,而且%replays也因输入而异。 - Zk1001
显示剩余7条评论
1个回答

3

谢谢提供参考,我会阅读它。顺便问一下,对y[]的访问是非常规律的,是吗?而且我只是写y[],根本没有读取它。另外,很抱歉在问题中可能没有提到所有的缺失都是缓存未命中读取。 - Zk1001

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接