CUDA найти максимальное значение в данном массиве

Question

Я пытался разработать небольшую программу CUDA для поиска максимального значения в данном массиве,

int input_data[0...50] = 1,2,3,4,5....,50

max_value инициализируется первым значением input_data[0], Окончательный ответ хранится в result[0]. Ядро дает 0 в качестве максимального значения. Я не знаю, в чем проблема. Я выполнил на 1 блоке 50 потоков.

__device__ int lock=0;

__global__ void max(float *input_data,float *result)
{
     float max_value = input_data[0];
     int  tid = threadIdx.x;

     if( input_data[tid] > max_value)
     {
         do{} while(atomicCAS(&lock,0,1));
         max_value=input_data[tid];
         __threadfence();
         lock=0;
      }

    __syncthreads();
    result[0]=max_value;  //Final result of max value 
}

Несмотря на то, что есть встроенные функции, я просто занимаюсь небольшими проблемами.

Answer 1

Вы пытаетесь создать «критическую секцию», но такой подход в CUDA может привести к зависанию всей вашей программы - старайтесь избегать ее, когда это возможно.

Почему ваш код зависает?

Ваше ядро ​​(__global__ функция) выполняется группами из 32 потоков, которые называются warps . Все потоки внутри одной деформации выполняются синхронно. Таким образом, деформация остановится в ваших do{} while(atomicCAS(&lock,0,1)), пока все нити из вашей деформации не получат блокировку. Но очевидно, что вы хотите предотвратить одновременное выполнение критической секции несколькими потоками. Это приводит к зависанию.

Альтернативное решение

Вам нужен «алгоритм параллельного сокращения». Вы можете начать читать здесь:

• Параллельный префикс sum @ wikipedia
• Параллельное сокращение @ веб-сайт CUDA
• Руководство NVIDIA по сокращению

Answer 2

Макс - это «сокращение» - посмотрите выборку Сокращения в SDK и сделайте максимум вместо суммирования.

Белая книга немного старая, но все еще достаточно полезная:

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

Последний шаг оптимизации заключается в использовании кодирования 'warp синхронный', чтобы избежать ненужных вызовов __syncthreads ().

Требуется как минимум 2 вызова ядра - один для записи группы промежуточных максимумовЗначения () для глобальной памяти, затем другое для определения max () этого массива.

Если вы хотите сделать это в одном вызове ядра, посмотрите пример SDK threadfenceReduction.Он использует __threadfence () и atomicAdd () для отслеживания прогресса, а затем 1 блок делает окончательное сокращение, когда все блоки закончили записывать свои промежуточные результаты.

Answer 3

Ваш код имеет потенциальную расу. Я не уверен, что вы определили переменную max_value в разделяемой памяти или нет, но оба неверны.

1) Если 'max_value' является просто локальной переменной, то каждый поток содержит ее локальную копию, которая не является фактическим максимальным значением (это просто максимальное значение между input_data [0] и input_data [tid]) , В последней строке кода все потоки записывают в результат [0] свое собственное значение max_value, что приведет к неопределенному поведению.

2) Если max_value - общая переменная, 49 потоков войдут в блок операторов if и попытаются обновить max_value по одному, используя блокировки. Но порядок выполнения среди 49 потоков не определен, и поэтому некоторые потоки могут перезаписывать фактическое максимальное значение на меньшие значения. Вам необходимо снова сравнить максимальное значение в критической секции.

Answer 4

Существуют разные способы доступа к переменным.когда вы определяете переменную с помощью device , тогда переменная помещается в глобальную память GPU и доступна всем потокам в сетке, shared помещает переменную в разделяемую память блока и становитсятолько потоками этого блока, в конце, если вы не используете какое-либо ключевое слово, например float max_value , переменная помещается в регистры потоков, и доступ к ней возможен только в этомthread.В вашем коде каждый поток имеет локальную переменную max_value, и он не идентифицирует переменные в других потоках.