Благодаря Grizzly у меня теперь работает ядро.Некоторые вещи, которые мне нужно было изменить, основаны на ответе Гризли:
Я добавил оператор IF в начале процедуры, чтобы отбросить все потоки, которые не будут ссылаться на какую-либо допустимую позицию в массивах, которые я использую.
if(featA > BLOCK_SIZE){return;}
При копировании первого дескриптора в локальную (разделяемую) память (от ig до B) необходимо указывать индекс, поскольку функция GetElement возвращает только один элемент за вызов (я пропустил этот вопрос в своем вопросе).
Bs[featA] = GetElement(tmpA, 0, featA);
Затем цикл SCAN нуждался в небольшой настройке, потому что буфер перезаписывался после каждой итерации, и никто не может контролировать, какой поток обращается к данным первым.Вот почему я «перерабатываю» буфер dif_acum для хранения частичных результатов и, таким образом, предотвращаю несогласованности в этом цикле.
dif_acum[featA] = accum[featA];
В цикле SCAN также есть некоторый граничный контроль для надежного определения условийдолжны быть добавлены вместе.
if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){
Наконец, я подумал, что имеет смысл включить приращение переменной цикла в последний оператор IF, чтобы его изменял только один поток.
if(featA == 0){
C[loop] = accum[BLOCK_SIZE-1];
loop += 1;
}
ЭтоЭто.Мне все еще интересно, как я могу использовать group_size, чтобы исключить это определение BLOCK_SIZE, и если есть более эффективные политики, которые я могу принять в отношении использования потоков.
Таким образом, код в конечном итоге выглядит так:
__kernel void CompareDescriptors(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE], __local float Bs[BLOCK_SIZE])
{
int gpidA = get_global_id(0);
int featA = get_local_id(0);
//global counter to store final differences
int loop = 0;
//auxiliary buffer to store temporary data
local float dif_acum[BLOCK_SIZE];
//discard the threads that are not going to be used.
if(featA > BLOCK_SIZE){
return;
}
//loop over all descriptors in A
for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){
//take the gpidA-th descriptor
DescriptorList tmpA = GetDescriptor(A, i);
//copy the current descriptor to local memory
Bs[featA] = GetElement(tmpA, 0, featA);
//loop over all the descriptors in B
for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
//take the difference of both current descriptors
dif_acum[featA] = Bs[featA]-B.elements[k*BLOCK_SIZE + featA];
//square the values in dif_acum
accum[featA] = dif_acum[featA]*dif_acum[featA];
barrier(CLK_LOCAL_MEM_FENCE);
//copy the values of accum to keep consistency once the scan procedure starts. Mostly important for the first element. Two buffers are necesarry because the scan procedure would override values that are then further read if one buffer is being used instead.
dif_acum[featA] = accum[featA];
//Compute the accumulated sum (a.k.a. scan)
for(int j = 1; j < BLOCK_SIZE; j *= 2){
int next_addend = featA-(j/2);
if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){
dif_acum[featA] = accum[featA] + accum[next_addend];
}
barrier(CLK_LOCAL_MEM_FENCE);
//copy As to accum
accum[featA] = GetElementArray(dif_acum, BLOCK_SIZE, featA);
barrier(CLK_LOCAL_MEM_FENCE);
}
//tell one of the threads to write the result of the scan in the array containing the results.
if(featA == 0){
C[loop] = accum[BLOCK_SIZE-1];
loop += 1;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
}
}