Варианты промахов кеша в GPU

Я использовал ядро ​​OpenCL, которое обращается к 7 глобальным буферам памяти, делает что-то со значениями и сохраняет результат обратно в 8-й глобальный буфер памяти. Как я заметил, по мере увеличения размера ввода коэффициент промахов кэша L1 (= промахи (промахи + попадания)) сильно меняется. Я не могу найти источник этой вариации. Размер ввода здесь означает количество глобальных рабочих элементов (степень двойки, кратная размеру рабочей группы). Число размеров рабочей группы остается 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.