Нахождение наименьшего числа в наборе с CUDA

Question

Нахождение наименьшего числа в наборе с CUDA

Предположим, у вас есть функция, которая будет принимать вектор, набор векторов и находить, какой вектор в наборе векторов был ближе всего к исходному вектору. Это может быть полезно, если я включил некоторый код:

int findBMU(float * inputVector, float * weights){int count = 0;
float currentDistance = 0;
int winner = 0;
float leastDistance = 99999;

for(int i = 0; i<10; i++){
for(int j = 0;j<10; j++){
for(int k = 0; k<10; k++){

int offset = (i*100+j*10+k)*644;
for(int i = offset; i<offset+644; i++){
currentDistance += abs((inputVector[count]-weights[i]))*abs((inputVector[count]-weights[i]));
count++;
}
currentDistance = sqrt(currentDistance);

count = 0;
if(currentDistance<leastDistance){
winner = offset;

leastDistance = currentDistance;

}
currentDistance = 0;
}
}
}
return winner;
}

В этом примере weights является одномерным массивом, с блоком из 644 элементов, соответствующих одному вектору. inputVector вектор, который сравнивается, и он также имеет 644 элемента.

Чтобы ускорить мою программу, я решил взглянуть на платформу CUDA, предоставленную NVIDIA. Вот как выглядел мой код, когда я изменил его в соответствии со спецификациями CUDA.

__global__ void findBMU(float * inputVector, float * weights, int * winner, float * leastDistance){int i = threadIdx.x+(blockIdx.x*blockDim.x);

if(i<1000){

int offset = i*644;
int count = 0;
float currentDistance = 0;
for(int w = offset; w<offset+644; w++){
currentDistance += abs((inputVector[count]-weights[w]))*abs((inputVector[count]-weights[w]));

count++;
}currentDistance = sqrt(currentDistance);

count = 0;
if(currentDistance<*leastDistance){
*winner = offset;

*leastDistance = currentDistance;

}
currentDistance = 0;
}

}

Для вызова функции я использовал: findBMU<<<20, 50>>>(d_data, d_weights, d_winner, d_least);

Но когда я вызываю функцию, иногда она дает мне правильный ответ, а иногда нет. Проведя некоторое исследование, я обнаружил, что у CUDA есть некоторые проблемы с такими проблемами сокращения, но я не мог найти, как это исправить. Как я могу изменить свою программу, чтобы она работала с CUDA?

0

c++cuda som vector

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

Проблема в том, что потоки, которые работают одновременно, будут видеть то же самое leastDistance и перезаписать результаты друг друга. Есть два значения, которые разделяются между потоками; leastDistance а также winner, У вас есть два основных варианта. Вы можете выписать результаты из всех потоков и затем сделать второй проход по данным с параллельным сокращением, чтобы определить, какой вектор имел лучшее соответствие, или вы можете реализовать это с помощью пользовательской атомарной операции, используя atomicCAS(),

Первый способ самый простой. Я предполагаю, что это также даст вам лучшую производительность, хотя и добавляет зависимость для бесплатной библиотеки Thrust. Вы бы использовали тяга :: min_element ().

Метод с использованием atomicCAS() использует тот факт, что atomicCAS() имеет 64-битный режим, в котором вы можете назначить любую семантику, которую вы хотите, 64-битному значению. В вашем случае вы бы использовали 32 бита для хранения leastDistance и 32 бита для хранения winner, Чтобы использовать этот метод, адаптируйте этот пример в Руководстве по программированию CUDA C, которое реализует плавающую точку двойной точности atomicAdd(),

__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull =
(unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
} while (assumed != old);
return __longlong_as_double(old);
}

1