Вы можете использовать clGetDeviceInfo, чтобы узнать, каков размер строки кэша для устройства.( clGetDeviceInfo , CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE) Сегодня на многих устройствах это значение обычно составляет 16 байт.
Маленькое чтение может быть проблематичным, но если вы читаете из той же строки кэша, у вас все будет хорошо.Краткий ответ: вам нужно держать свои «маленькие кусочки» близко друг к другу в памяти, чтобы они работали быстро.
У меня есть две функции ниже, чтобы продемонстрировать два способа доступа к памяти - vectorAddFoo и vectorAddBar.Третья функция copySomeMemory (...) относится именно к вашему вопросу.Обе функции вектора имеют свои рабочие элементы, которые добавляют часть добавляемых векторов, но используют разные шаблоны доступа к памяти.vectorAddFoo заставляет каждый рабочий элемент обрабатывать блок векторных элементов, начиная с его вычисленной позиции в массивах и продвигаясь вперед через свою рабочую нагрузку.В vectorAddBar рабочие элементы начинаются с их gid и пропускают элементы gSize (= глобальный размер) перед извлечением и добавлением следующих элементов.
vectorAddBar будет выполняться быстрее из-за операций чтения и записи, попадающих в одну и ту же кэш-строку в памяти.Каждые 4 чтения с плавающей запятой будут попадать в одну и ту же строку кэша и выполнять только одно действие из контроллера памяти для выполнения.После прочтения a [] и b [] в этом вопросе все четыре рабочих элемента смогут выполнить свое добавление и поставить в очередь свою запись в c [].
vectorAddFoo гарантирует, что операции чтения и записи не находятся вта же самая кешлайн (за исключением очень коротких векторов ~ totalElements <5).Каждое чтение из рабочего элемента потребует действия контроллера памяти.Если в каждом случае gpu не кэширует следующие 3 числа с плавающей запятой, это приведет к четырехкратному доступу к памяти. </p>
__kernel void
vectorAddFoo(__global const float * a,
__global const float * b,
__global float * c,
__global const totalElements)
{
int gid = get_global_id(0);
int elementsPerWorkItem = totalElements/get_global_size(0);
int start = elementsPerWorkItem * gid;
for(int i=0;i<elementsPerWorkItem;i++){
c[start+i] = a[start+i] + b[start+i];
}
}
__kernel void
vectorAddBar(__global const float * a,
__global const float * b,
__global float * c,
__global const totalElements)
{
int gid = get_global_id(0);
int gSize = get_global_size(0);
for(int i=gid;i<totalElements;i+=gSize){
c[i] = a[i] + b[i];
}
}
__kernel void
copySomeMemory(__global const int * src,
__global const count,
__global const position)
{
//copy 16kb of integers to local memory, starting at 'position'
int start = position + get_local_id(0);
int lSize = get_local_size(0);
__local dst[4096];
for(int i=0;i<4096;i+=lSize ){
dst[start+i] = src[start+i];
}
barrier(CLK_GLOBAL_MEM_FENCE);
//use dst here...
}