Изменение пропусков кеша в GPU

Я собирал ядро 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.

Ответ 1