Я играю в ядре 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.