Этот пример показывает, как простая, известная математическая проблема, Множество Мандельброта, может быть выражена в коде MATLAB®. Используя Parallel Computing Toolbox™ этот код затем адаптируется, чтобы использовать оборудование графического процессора тремя способами:
Используя существующий алгоритм, но с данными графического процессора, как введено
Используя arrayfun
, чтобы выполнить алгоритм на каждом элементе независимо
Используя интерфейс MATLAB/CUDA, чтобы запустить некоторый существующий CUDA/C ++ код
Значения ниже задают высоко масштабируемую часть Множества Мандельброта в долине между основной кардиоидой и p/q лампой с ее левой стороны от него.
1000x1000 сетка действительных частей (X) и мнимых частей (Y) создается между этими пределами, и алгоритм Мандельброта выполнен с помощью итераций в каждом местоположении сетки. Для этого конкретного местоположения 500 итераций будет достаточно, чтобы полностью представить изображение.
maxIterations = 500; gridSize = 1000; xlim = [-0.748766713922161, -0.748766707771757]; ylim = [ 0.123640844894862, 0.123640851045266];
Ниже реализация Множества Мандельброта с помощью стандартных команд 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
. Для входных параметров графического процессора массивов функция, используемая с arrayfun
, скомпилирована в нативный код графического процессора. В этом случае мы поместили цикл в pctdemo_processMandelbrotElement.m
:
function count = (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);
Обратите внимание на то, что раннее аварийное прекращение работы было введено потому что этот функциональные процессы только один элемент. Для большинства представлений Множества Мандельброта значительное количество элементов останавливается очень рано, и это может сохранить большую обработку. Цикл for
был также заменен циклом while
, потому что они обычно более эффективны. Эта функция не упоминает о графическом процессоре и не использует специфичных для графического процессора функций - это - стандартный код MATLAB.
Используя средние значения arrayfun
, что вместо многих тысяч вызовов, чтобы разделить оптимизированные графическим процессором операции (по крайней мере 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 ) );
В Экспериментах в MATLAB улучшенная производительность достигается путем преобразования основного алгоритма для C-MEX-функции. Если вы готовы сделать, некоторые работают на C/C++, то можно использовать Parallel Computing Toolbox, чтобы вызвать предзаписанное использование ядер CUDA данные MATLAB. Вы делаете это с функцией parallel.gpu.CUDAKernel
.
CUDA/C ++ реализация алгоритма обработки элемента был написан от руки в pctdemo_processMandelbrotElement.cu
: Это должно затем быть вручную скомпилировано с помощью компилятора Nvidia NVCC, чтобы произвести уровень ассемблера pctdemo_processMandelbrotElement.ptx
(.ptx
обозначает "Параллельный язык выполнения Потока").
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; }
Один поток графического процессора требуется для местоположения во Множестве Мандельброта с потоками, сгруппированными в блоки. Ядро указывает, насколько большой блок потока, и в коде ниже мы используем это, чтобы вычислить количество требуемых блоков потока. Это затем становится 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 может быть адаптирован, чтобы использовать оборудование графического процессора:
Преобразуйте входные данные, чтобы быть на графическом процессоре с помощью gpuArray
, оставив алгоритм без изменений
Используйте arrayfun
на входе gpuArray
, чтобы выполнить алгоритм на каждом элементе входа независимо
Используйте parallel.gpu.CUDAKernel
, чтобы запустить некоторый существующий CUDA/C ++ использование кода данные MATLAB
title('The Mandelbrot Set on a GPU')