Sie können clGetDeviceInfo verwenden, um herauszufinden, wie groß die Cacheline für ein Gerät ist. (clGetDeviceInfo, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE) Dieser Wert ist heute auf vielen Geräten normalerweise 16 Byte.
Kleine Lesevorgänge können mühsam sein, aber wenn Sie von der gleichen Cachezeile lesen, sollten Sie in Ordnung sein. Die kurze Antwort: Sie müssen Ihre "kleinen Brocken" im Speicher nahe beieinander halten, um sie schnell zu halten.
Ich habe zwei Funktionen unten, um zwei Möglichkeiten zu zeigen, auf den Speicher zuzugreifen - vectorAddFoo und vectorAddBar. Die dritte Funktion copySomeMemory (...) bezieht sich speziell auf Ihre Frage. Beide Vektorfunktionen fügen ihre Arbeitsaufgaben einen Teil der hinzugefügten Vektoren hinzu, verwenden jedoch unterschiedliche Speicherzugriffsmuster. vectorAddFoo ruft jedes Arbeitselement auf, um einen Block von Vektorelementen zu verarbeiten, beginnend bei seiner berechneten Position in den Arrays und vorwärts durch seine Arbeitslast. vectorAddBar hat Arbeitselemente, die an ihrem gid beginnen und gSize-Elemente (= globale Größe) überspringen, bevor sie die nächsten Elemente abrufen und hinzufügen.
vectorAddBar wird schneller ausgeführt, da die Lese- und Schreibvorgänge in dieselbe Cachezeile im Speicher fallen. Alle 4 Float-Lesevorgänge werden auf die gleiche Cache-Line fallen und nur eine Aktion vom Memory-Controller ausführen. Nach dem Lesen von a [] und b [] in dieser Angelegenheit können alle vier Arbeitselemente ihre Addition durchführen und ihre Schreibvorgänge in c [] einreihen.
vectorAddFoo wird garantieren, dass die Lese- und Schreibvorgänge nicht in der gleichen Cachezeile sind (außer für sehr kurze Vektoren ~ totalElements < 5). Jedes Lesen von einem Arbeitselement erfordert eine Aktion vom Speichercontroller. Sofern die GPU nicht die folgenden 3 Floats in jedem Fall zwischenspeichert, führt dies zu einem 4x Speicherzugriff.
__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...
}