Установка

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

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

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

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

Ниже приведена реализация набора Mandelbrot с использованием стандартных команд MATLAB, выполняемых на CPU. Это основано на коде, предоставленном в электронной книге Клеве Молер «Эксперименты с 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 обнаруживает данные на GPU, вычисления с этими данными выполняются на GPU. Класс gpuArray предоставляет версии графических процессоров многих функций, которые можно использовать для создания массивов данных, включая linspace, logspace, и meshgrid здесь необходимы функции. Аналогично, count инициализируется непосредственно на GPU с помощью функции 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 компилируется в собственный код графического процессора. В этом случае мы поместили цикл в 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);

Следует отметить, что было введено раннее прекращение, поскольку эта функция обрабатывает только один элемент. Для большинства видов набора Mandelbrot значительное количество элементов останавливается очень рано и это может сэкономить много обработки. for цикл также был заменен на while цикл, потому что они, как правило, более эффективны. Эта функция не упоминает GPU и не использует специфичные для 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 означает «язык параллельного потока 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 требуется один поток графического процессора, сгруппированный в блоки. Ядро указывает, насколько велик блок потока, и в приведенном ниже коде мы используем его для вычисления требуемого количества блоков потока. Это затем становится 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')