Ядра из Рассеяния - операции сбора типов

GPU Coder™ также поддерживает концепцию сокращений - важное исключение из правила, что итерации цикла должны быть независимыми. Переменная сокращения накапливает значение, которое зависит от всех итераций вместе, но не зависит от порядка итерации. Переменные сокращения появляются на обеих сторонах оператора назначения, таких как суммирование, скалярный продукт и сортировка. В следующем примере показано типовое использование переменной сокращения x:

x = ...; % Some initialization of x
for i = 1:n
  x = x + d(i);
end

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

В древовидном подходе можно выполнять каждый горизонтальный уровень дерева параллельно в течение определенного количества проходов. По сравнению с последовательным выполнением двоичное дерево требует большей памяти, потому что каждый проход требует массива временных значений в качестве вывода. Преимущества эффективности, которые вы получаете, значительно перевешивают стоимость увеличения использования памяти. GPU Coder создает ядра восстановления с помощью этого древовидного подхода, в котором каждый блок потока уменьшает фрагмент массива. Параллельное сокращение требует частичного обмена данными результатов между блоками потоков. В более старых CUDA^® устройства, этот обмен данными был достигнут за счет использования общей памяти и синхронизации потоков. Начиная с архитектуры GPU Kepler, CUDA обеспечивает shuffle (shfl) инструкция и быстрые атомарные операции памяти устройства, которые делают сокращения еще быстрее. Уменьшайте ядра, которые создает GPU Coder, используйте shfl_down инструкция для уменьшения на основе (32 потока) потоков. Затем первый поток каждой основы использует инструкции по атомарной операции, чтобы обновить уменьшенное значение.

Для получения дополнительной информации о инструкциях см. NVIDIA^® документация.

Пример векторной суммы

В этом примере показано, как создать ядра типов сокращения CUDA с помощью GPU Coder. Предположим, что вы хотите создать вектор v и вычислите сумму его элементов. Можно реализовать этот пример как MATLAB^® функция.

function s = VecSum(v)
    s = 0;
    for i = 1:length(v)
       s = s + v(i);
    end
end

Подготовьте vecSum для создания ядра

GPU Coder не требует специальной прагмы для вывода восстановительных ядер. В этом примере используйте coder.gpu.kernelfun прагма для генерации восстановительных ядер CUDA. Используйте измененную VecSum функция.

function s = VecSum(v) %#codegen
    s = 0;
    
    coder.gpu.kernelfun();
    for i = 1:length(v)
       s = s + v(i);
    end
end

Сгенерированный код CUDA

Когда вы генерируете код CUDA с помощью приложения GPU Coder или из командной строки, GPU Coder создает одно ядро, которое выполняет вычисление суммы векторов. Ниже представлен фрагмент vecSum_kernel1.

static __global__ __launch_bounds__(512, 1) void vecSum_kernel1(const real_T *v,
  real_T *s)
{
  uint32_T threadId;
  uint32_T threadStride;
  uint32_T thdBlkId;
  uint32_T idx;
  real_T tmpRed;
  ;
  ;
  thdBlkId = (threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x)
    + threadIdx.x;
  threadId = ((gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y) +
              blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) + thdBlkId;
  threadStride = gridDim.x * blockDim.x * (gridDim.y * blockDim.y) * (gridDim.z *
    blockDim.z);
  if (!((int32_T)threadId >= 512)) {
    tmpRed = 0.0;
    for (idx = threadId; threadStride < 0U ? idx >= 511U : idx <= 511U; idx +=
         threadStride) {
      tmpRed += v[idx];
    }

    tmpRed = workGroupReduction1(tmpRed, 0.0);
    if (thdBlkId == 0U) {
      atomicOp1(s, tmpRed);
    }
  }
}

Перед вызовом VecSum_kernel1, два cudaMemcpy вызывает передачу вектора v и скалярная s от хоста к устройству. Ядро имеет один блок потока, содержащий 512 потоков на блок, согласующийся с размером входного вектора. Третий cudaMemcpy вызов копирует результат расчета обратно на хост. Ниже приведен фрагмент основной функции.

  cudaMemcpy((void *)gpu_v, (void *)v, 4096ULL, cudaMemcpyHostToDevice);
  cudaMemcpy((void *)gpu_s, (void *)&s, 8ULL, cudaMemcpyHostToDevice);
  VecSum_kernel1<<<dim3(1U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_v, gpu_s);
  cudaMemcpy(&s, gpu_s, 8U, cudaMemcpyDeviceToHost);

Примечание

Для лучшей эффективности GPU Coder отдает приоритет параллельным ядрам по сравнению с сокращениями. Если ваш алгоритм содержит сокращение внутри параллельного цикла, GPU Coder выводит сокращение как регулярный цикл и генерирует для него ядра.

1-D операций сокращения на графическом процессоре

Вы можете использовать gpucoder.reduce функция для генерации кода CUDA, который выполняет эффективные операции сокращения 1-D на графическом процессоре. Сгенерированный код использует признаки shuffle CUDA для реализации операции сокращения.

Например, чтобы найти sum и max элементы массива A:

function s = myReduce(A)
   s = gpucoder.reduce(A, {@mysum, @mymax}); 
end

function c = mysum(a, b)
   c = a+b;
end

function c = mymax(a, b)
   c = max(a,b);
end

Для генерации кода, gpucoder.reduce функция имеет следующие требования:

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

Примечание

Для некоторых входов целочисленного типа данных код, сгенерированный для gpucoder.reduce функция может содержать промежуточные расчеты, которые достигают насыщения. В таких случаях результаты из сгенерированного кода могут не совпадать с результатами симуляции из MATLAB.

См. также

Документация