Setup

Приведенные ниже значения определяют сильно увеличенную часть набора Мандельброта в овраге между основным кардиоидом и p/q луковицей налево.

Между этими пределами создается сетка 1000x1000 из действительных частей (X) и мнимых частей (Y), и алгоритм Мандельброта итератируется в каждом месте сетки. Для этого конкретного местоположения будет достаточно 500 итераций, чтобы полностью отобразить изображение.

maxIterations = 500;
gridSize = 1000;
xlim = [-0.748766713922161, -0.748766707771757];
ylim = [ 0.123640844894862,  0.123640851045266];

Набор Мандельброта в MATLAB

Ниже представлена реализация набора Mandelbrot Set с использованием стандартных команд MATLAB, выполняемых на центральном процессоре. Это основано на коде, приведенном в электронной книге Клева Молера «Эксперименты с MATLAB».

Это вычисление векторизировано таким образом, что каждое местоположение обновляется сразу.

% Setup
t = tic();
x = linspace( xlim(1), xlim(2), gridSize );
y = linspace( ylim(1), ylim(2), gridSize );
[xGrid,yGrid] = meshgrid( x, y );
z0 = xGrid + 1i*yGrid;
count = ones( size(z0) );

% Calculate
z = z0;
for n = 0:maxIterations
    z = z.*z + z0;
    inside = abs( z )<=2;
    count = count + inside;
end
count = log( count );

% Show
cpuTime = toc( t );
fig = gcf;
fig.Position = [200 200 600 600];
imagesc( x, y, count );
colormap( [jet();flipud( jet() );0 0 0] );
axis off
title( sprintf( '%1.2fsecs (without GPU)', cpuTime ) );

Использование gpuArray

Когда MATLAB встречается с данными о графическом процессоре, вычисления с этими данными выполняются на графическом процессоре. Класс gpuArray предоставляет версии графический процессор многих функций, которые можно использовать для создания массивов данных, включая linspace, logspace, и meshgrid функции, необходимые здесь. Точно так же count массив инициализируется непосредственно на графическом процессоре с помощью функции ones.

С этими изменениями в инициализации данных вычисления теперь будут выполняться на графическом процессоре:

% Setup
t = tic();
x = gpuArray.linspace( xlim(1), xlim(2), gridSize );
y = gpuArray.linspace( ylim(1), ylim(2), gridSize );
[xGrid,yGrid] = meshgrid( x, y );
z0 = complex( xGrid, yGrid );
count = ones( size(z0), 'gpuArray' );

% Calculate
z = z0;
for n = 0:maxIterations
    z = z.*z + z0;
    inside = abs( z )<=2;
    count = count + inside;
end
count = log( count );

% Show
count = gather( count ); % Fetch the data back from the GPU
naiveGPUTime = toc( t );
imagesc( x, y, count )
axis off
title( sprintf( '%1.3fsecs (naive GPU) = %1.1fx faster', ...
    naiveGPUTime, cpuTime/naiveGPUTime ) )

Поэлементная операция

Отметив, что алгоритм работает одинаково на каждом элементе входа, мы можем поместить код в вспомогательную функцию и вызвать его с помощью arrayfun. Для входов массива GPU, функция, используемая с arrayfun компилируется в собственный код GPU. В этом случае мы поместили цикл в pctdemo_processMandelbrotElement.m:

function count = pctdemo_processMandelbrotElement(x0,y0,maxIterations)
z0 = complex(x0,y0);
z = z0;
count = 1;
while (count <= maxIterations) && (abs(z) <= 2)
    count = count + 1;
    z = z*z + z0;
end
count = log(count);

Обратите внимание, что раннее прекращение было введено, потому что эта функция обрабатывает только один элемент. Для большинства представлений набора Мандельброта значительное количество элементов останавливается очень рано, и это может сэкономить много обработки. The for цикл также был заменен на while цикл, потому что они обычно более эффективны. Эта функция не упоминает о графическом процессоре и не использует специфичных для GPU функций - это стандартный код MATLAB.

Использование arrayfun означает, что вместо многих тысяч вызовов для разделения оптимизированных для GPU операций (не менее 6 на итерацию) мы делаем один вызов на параллельную операцию GPU, которая выполняет весь расчет. Это значительно уменьшает накладные расходы.

% Setup
t = tic();
x = gpuArray.linspace( xlim(1), xlim(2), gridSize );
y = gpuArray.linspace( ylim(1), ylim(2), gridSize );
[xGrid,yGrid] = meshgrid( x, y );

% Calculate
count = arrayfun( @pctdemo_processMandelbrotElement, ...
                  xGrid, yGrid, maxIterations );

% Show
count = gather( count ); % Fetch the data back from the GPU
gpuArrayfunTime = toc( t );
imagesc( x, y, count )
axis off
title( sprintf( '%1.3fsecs (GPU arrayfun) = %1.1fx faster', ...
    gpuArrayfunTime, cpuTime/gpuArrayfunTime ) );

Работа с CUDA

В Экспериментах в MATLAB улучшенная производительность достигается путем преобразования базового алгоритма в функцию C-Mex. Если вы готовы выполнить некоторую работу в C/C + +, то можно использовать Parallel Computing Toolbox, чтобы вызвать предварительно написанные ядра CUDA с помощью данных MATLAB. Вы делаете это с parallel.gpu.CUDAKernel функция.

Реализация CUDA/C + + алгоритма обработки элемента написана вручную pctdemo_processMandelbrotElement.cu: Затем это должно быть скомпилировано вручную с помощью компилятора NVCC от nVidia, чтобы создать pctdemo_processMandelbrotElement.ptx уровня сборки (.ptx означает «Язык Parallel Thread eXecution»).

Код CUDA/C + + чуть более вовлечен, чем версии MATLAB, которые мы видели до сих пор, из-за отсутствия сложных чисел в C++. Однако суть алгоритма неизменна:

__device__
unsigned int doIterations( double const realPart0,
                           double const imagPart0,
                           unsigned int const maxIters ) {
   // Initialize: z = z0
   double realPart = realPart0;
   double imagPart = imagPart0;
   unsigned int count = 0;
   // Loop until escape
   while ( ( count <= maxIters )
          && ((realPart*realPart + imagPart*imagPart) <= 4.0) ) {
      ++count;
      // Update: z = z*z + z0;
      double const oldRealPart = realPart;
      realPart = realPart*realPart - imagPart*imagPart + realPart0;
      imagPart = 2.0*oldRealPart*imagPart + imagPart0;
   }
   return count;
}

Для расположения в наборе Mandelbrot требуется один поток графический процессор, причем потоки сгруппированы в блоки. Ядро указывает, насколько велик блок thread-block, и в коде ниже мы используем это, чтобы вычислить количество необходимых блоков thread-block. Это затем становится GridSize.

% Load the kernel
cudaFilename = 'pctdemo_processMandelbrotElement.cu';
ptxFilename = ['pctdemo_processMandelbrotElement.',parallel.gpu.ptxext];
kernel = parallel.gpu.CUDAKernel( ptxFilename, cudaFilename );

% Setup
t = tic();
x = gpuArray.linspace( xlim(1), xlim(2), gridSize );
y = gpuArray.linspace( ylim(1), ylim(2), gridSize );
[xGrid,yGrid] = meshgrid( x, y );

% Make sure we have sufficient blocks to cover all of the locations
numElements = numel( xGrid );
kernel.ThreadBlockSize = [kernel.MaxThreadsPerBlock,1,1];
kernel.GridSize = [ceil(numElements/kernel.MaxThreadsPerBlock),1];

% Call the kernel
count = zeros( size(xGrid), 'gpuArray' );
count = feval( kernel, count, xGrid, yGrid, maxIterations, numElements );

% Show
count = gather( count ); % Fetch the data back from the GPU
gpuCUDAKernelTime = toc( t );
imagesc( x, y, count )
axis off
title( sprintf( '%1.3fsecs (GPU CUDAKernel) = %1.1fx faster', ...
    gpuCUDAKernelTime, cpuTime/gpuCUDAKernelTime ) );

Сводные данные

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

title('The Mandelbrot Set on a GPU')