CUDA сокращение с помощью регистров - PullRequest
Новая
       47

CUDA сокращение с помощью регистров

1 голос
/ 22 июля 2011

Мне нужно рассчитать средние значения N сигналов, используя уменьшение.Вход представляет собой одномерный массив размера MN, где M - длина каждого сигнала.

Первоначально у меня была дополнительная общая память, чтобы сначала скопировать данные и выполнить сокращение для каждого сигнала.Однако исходные данные повреждены.

Моя программа пытается минимизировать общую память.Поэтому мне было интересно, как я могу использовать регистры для суммирования N сигналов.У меня N потоков, разделяемая память (float) s_m [N * M], 0 .... M-1 - первый сигнал и т. Д.

Нужно ли мне хранить N регистров (или один)среднее значение N разных сигналов?(Я знаю, как сделать с последовательным добавлением, используя многопоточное программирование и 1 регистр).Следующий шаг, который я хочу сделать, - вычесть каждое значение на входе из среднего значения соответствующего ему сигнала.

Ответы [ 2 ]

3 голосов
/ 10 сентября 2012

Ваша проблема очень мала (N = 32 и M <128).Однако некоторые рекомендации: </p>

Предполагается, что вы сокращаете значения N для каждого из N потоков.

  • Если N очень большое (> 10 тысяч), просто сделайте сокращенияболее M последовательно в каждом потоке.
  • Если N <10 с тысяч, рассмотрите возможность использования одного деформации или одного блока потока для выполнения каждого из N сокращений. </li>
  • Если N очень мало, но Mочень большой, рассмотрите возможность использования нескольких потоковых блоков для каждого из сокращений N.
  • Если N очень мало, а M очень мало (как и ваши числа), рассмотрите возможность использования только графического процессора для сокращений, если вычислениякоторые генерируют и / или потребляют ввод / вывод сокращений, также работают на GPU.
0 голосов
/ 24 июля 2011

Исходя из моего понимания вопроса, я говорю, что вам не нужны N регистров для хранения среднего значения N различных сигналов.

Если у вас уже есть N потоков [Учитывая, что каждый поток выполняет сокращение только для одного сигнала], то вам не нужны N регистров для хранения сокращения одного сигнала. Все, что вам нужно, это один регистр для хранения среднего значения. 

dim3 threads (N,1);
reduction<<<threads,1>>>(signals);  // signals is the [N*M] array 

__global__ reduction (int *signals)
{
   int id = threadIdx.x;
   float meanValue = 0.0;

   for(int i = 0; i < M; i++)
          meanValue = signals[id*M +i];

   meanValue =  meanValue/M;

   // Then do the subtraction
   for(int i = 0; i < M; i++)
          signals[id*M +i] -= meanValue;
}

Если вам нужно выполнить глобальное уменьшение всех средних значений N различных сигналов, то вам нужно использовать 2 регистра [один для хранения локального среднего и другой для хранения глобального среднего] и разделяемой памяти 

dim3 threads (N,1);
reduction<<<threads,1>>>(signals);  // signals is the [N*M] array 

__global__ reduction (int *signals)
{
   __shared__ float means[N];      // shared value
   int id = threadIdx.x;
   float meanValue = 0.0;
   float globalMean = 0.0;

   for(int i = 0; i < M; i++)
          meanValue += signals[id*M +i];

   means[id] =  meanValue/M;

   __syncthreads();

   // do the global reduction
   for(int i = 0; i < N; i++)
          globalMean += means[i];

   globalMean = globalMean/N;

   // Then do the subtraction
   for(int i = 0; i < M; i++)
          signals[id*M +i] -= globalMean;
}

Надеюсь, это поможет вам. Любые сомнения, дайте мне знать.

...