Эта тема объясняет, как создать исполняемое ядро из CU или PTX (параллельное выполнение потока) файлы и запуск что ядро на графическом процессоре от MATLAB®. Ядро представлено в MATLAB CUDAKernel
объект, который может работать с массивом MATLAB или gpuArray переменными.
Следующие шаги описывают общий рабочий процесс CUDAKernel:
Используйте скомпилированный код PTX, чтобы создать объект CUDAKernel, который содержит исполняемый код графического процессора.
Установите свойства на объекте CUDAKernel управлять его выполнением на графическом процессоре.
Вызовите feval
на CUDAKernel с необходимыми входными параметрами, чтобы запустить ядро на графическом процессоре.
Код MATLAB, который выполняет эти шаги, может выглядеть примерно так:
% 1. Create CUDAKernel object. k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','entryPt1'); % 2. Set object properties. k.GridSize = [8 1]; k.ThreadBlockSize = [16 1]; % 3. Call feval with defined inputs. g1 = gpuArray(in1); % Input gpuArray. g2 = gpuArray(in2); % Input gpuArray. result = feval(k,g1,g2);
Следующие разделы предоставляют подробную информацию этих команд и шагов рабочего процесса.
Если у вас есть файл CU, вы хотите выполниться на графическом процессоре, необходимо сначала скомпилировать его, чтобы создать файл PTX. Один способ сделать это с nvcc
компилятор в NVIDIA® CUDA® Toolkit. Например, если ваш файл CU называется myfun.cu
, можно создать скомпилированный файл PTX с командой интерпретатора:
nvcc -ptx myfun.cu
Это генерирует файл с именем myfun.ptx
.
С .cu
файл и .ptx
файл можно создать CUDAKernel
объект в MATLAB, который можно затем использовать, чтобы оценить ядро:
k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');
Вы не можете save
или load
Объекты CUDAKernel.
Если у вас нет файла CU, соответствующего вашему файлу PTX, можно задать прототип C для ядра C вместо файла CU. Например:
k = parallel.gpu.CUDAKernel('myfun.ptx','float *, const float *, float');
Другое использование для прототипного входа C состоит в том, когда ваш исходный код использует нераспознанное переименование поддерживаемого типа данных. (См. поддерживаемые типы ниже.) Предположим ваше ядро включает следующий код.
typedef float ArgType; __global__ void add3( ArgType * v1, const ArgType * v2 ) { int idx = threadIdx.x; v1[idx] += v2[idx]; }
ArgType
самостоятельно не распознан поддерживаемым типом данных, таким образом, файл CU, который включает его, не может непосредственно использоваться в качестве входа при создании объекта CUDAKernel в MATLAB. Однако поддерживаемый вход вводит к add3
ядро может быть задано как C прототипный вход конструктору CUDAKernel. Например:
k = parallel.gpu.CUDAKernel('test.ptx','float *, const float *','add3');
Поддерживаемые типы данных стандарта C/C++ перечислены в следующей таблице.
Типы плавающие | Целочисленные типы | Булевы и типы символов |
---|---|---|
|
|
|
Кроме того, следующие целочисленные типы поддерживаются, когда вы включаете tmwtypes.h
заголовочный файл в вашей программе.
Целочисленные типы |
---|
|
Заголовочный файл поставляется как
. Вы включаете файл в свою программу с линией:matlabroot
/extern/include/tmwtypes.h
#include "tmwtypes.h"
Все входные параметры могут быть скалярами или указателями, и могут быть помечены const
.
Объявление C ядра всегда имеет форму:
__global__ void aKernel(inputs ...)
Ядро ничего не должно возвращать и работать только с его входными параметрами (скаляры или указатели).
Ядро не может выделить любую форму памяти, таким образом, все выходные параметры должны быть предварительно выделены, прежде чем ядро выполняется. Поэтому размеры всех выходных параметров должны быть известны, прежде чем вы запустите ядро.
В принципе все указатели передали в ядро, которые не являются const
мог содержать выходные данные, поскольку много потоков ядра могли изменить те данные.
При переводе определения ядра в C в MATLAB:
Весь скаляр вводит в C (double
, float
, int
, и т.д.), должны быть скаляры в MATLAB или скаляр (т.е. одноэлементный) gpuArray переменные.
Весь const
указатель вводит в C (const double *
, и т.д.), могут быть скаляры или матрицы в MATLAB. Они брошены к правильному типу, скопированному на устройство, и указатель на первый элемент передается ядру. Никакая информация о первоначальном размере не передается ядру. Это - как будто ядро непосредственно получило результат mxGetData
на mxArray
.
Все непостоянные входные параметры указателя в C передаются ядру точно как непостоянные указатели. Однако, потому что непостоянный указатель мог быть изменен ядром, это будет рассмотрено как выход от ядра.
Входные параметры от скаляров рабочего пространства MATLAB и массивов брошены в требуемый тип и затем переданы ядру. Однако входные параметры gpuArray автоматически не брошены, таким образом, их тип и сложность должны точно совпадать с ожидаемыми.
Эти правила имеют некоторые последствия. Самое известное - то, что каждый выход от ядра должен обязательно также быть входом к ядру, поскольку вход позволяет пользователю задавать размер выхода (который следует из неспособности выделить память на графическом процессоре).
Когда вы создаете объект ядра без завершающей работу точки с запятой, или когда вы вводите переменную объекта в командной строке, MATLAB отображает свойства объекта ядра. Например:
k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu')
k = parallel.gpu.CUDAKernel handle Package: parallel.gpu Properties: ThreadBlockSize: [1 1 1] MaxThreadsPerBlock: 512 GridSize: [1 1 1] SharedMemorySize: 0 EntryPoint: '_Z8theEntryPf' MaxNumLHSArguments: 1 NumRHSArguments: 2 ArgumentTypes: {'in single vector' 'inout single vector'}
Свойства объекта ядра управляют частью его поведения при выполнении. Используйте запись через точку, чтобы изменить те свойства, которые могут быть изменены.
Для описания свойств объектов, смотрите CUDAKernel
страница ссылки на объект. Типичная причина изменения устанавливаемых свойств состоит в том, чтобы задать количество потоков, аналогичных описанному ниже.
Если ваш файл PTX содержит многократные точки, можно идентифицировать конкретное ядро в myfun.ptx
то, что вы хотите объект ядра k
относиться к:
k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','myKernel1');
Один файл PTX может содержать многократные точки к различным ядрам. Каждая из этих точек входа имеет уникальное имя. Эти имена обычно искажаются (как в искажении C++). Однако, когда сгенерировано nvcc
имя PTX всегда содержит исходное имя функции из файла CU. Например, если файл CU задает функцию ядра как
__global__ void simplestKernelEver( float * x, float val )
затем код PTX содержит запись, которая может быть названа _Z18simplestKernelEverPff
.
Когда вы имеете многократные точки, задаете имя записи для конкретного ядра при вызове CUDAKernel
сгенерировать ваше ядро.
CUDAKernel
функционируйте ищет ваше имя записи в файле PTX и соответствует на любых случаях подстроки. Поэтому вы не должны называть ни одну из своих записей как подстроки никаких других.
Вы не можете управлять исходными именами записи, в этом случае необходимо знать об уникальном, искаженном выведенный для каждого. Например, рассмотрите следующий шаблон функции.
template <typename T> __global__ void add4( T * v1, const T * v2 ) { int idx = threadIdx.x; v1[idx] += v2[idx]; }
Когда шаблон расширен для плавания и дважды, это приводит к двум точкам входа, обе из которых содержат подстроку add4
.
template __global__ void add4<float>(float *, const float *); template __global__ void add4<double>(double *, const double *);
PTX имеет соответствующие записи:
_Z4add4IfEvPT_PKS0_ _Z4add4IdEvPT_PKS0_
Используйте точку входа add4If
для версии плавающей и add4Id
для двойной версии.
k = parallel.gpu.CUDAKernel('test.ptx','double *, const double *','add4Id');
Вы задаете количество вычислительных потоков для вашего CUDAKernel путем установки двух из его свойств объектов:
GridSize
— Вектор трех элементов, продукт которых определяет количество блоков.
ThreadBlockSize
— Вектор трех элементов, продукт которых определяет количество потоков на блок. (Обратите внимание на то, что продукт не может превысить значение свойства MaxThreadsPerBlock
.)
Значением по умолчанию для обоих из этих свойств является [1 1 1]
, но предположите, что вы хотите использовать 500 потоков, чтобы запустить поэлементные операции на векторах 500 элементов параллельно. Простой способ достигнуть этого состоит в том, чтобы создать ваш CUDAKernel и установить его свойства соответственно:
k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu'); k.ThreadBlockSize = [500,1,1];
Обычно вы устанавливаете сетку и распараллеливаете размеры блока на основе размеров ваших входных параметров. Для получения информации об иерархии потока, и сетках нескольких-размерностей и блоках, смотрите NVIDIA CUDA C Руководство по программированию.
Используйте feval
функция, чтобы оценить CUDAKernel на графическом процессоре. Следующие примеры показывают, как выполнить ядро с помощью переменных рабочего пространства MATLAB и gpuArray переменных.
Примите, что вы уже написали некоторые ядра в родном языке и хотите использовать их в MATLAB, чтобы выполниться на графическом процессоре. У вас есть ядро, которое делает свертку на двух векторах; загрузите и запустите его с двумя случайными входными векторами:
k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu'); result = feval(k,rand(100,1),rand(100,1));
Даже если входные параметры являются константами или переменными для данных о рабочем пространстве MATLAB, выходом является gpuArray
.
Может быть более эффективно использовать gpuArray
объекты, как введено при выполнении ядра:
k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu'); i1 = gpuArray(rand(100,1,'single')); i2 = gpuArray(rand(100,1,'single')); result1 = feval(k,i1,i2);
Поскольку выходом является gpuArray
, можно теперь выполнить другие операции с помощью этих данных о вводе или выводе без дальнейших передач между рабочим пространством MATLAB и графическим процессором. Когда все ваши расчеты графического процессора завершены, собирают ваши данные о конечном результате в рабочее пространство MATLAB:
result2 = feval(k,i1,i2); r1 = gather(result1); r2 = gather(result2);
При вызове [out1, out2] = feval(kernel, in1, in2, in3)
, входные параметры in1
, in2
, и in3
соответствуйте каждому из входных параметров к функции C в вашем файле CU. Выходные параметры out1
и out2
сохраните значения первых и вторых входных параметров указателя неconst к функции C после того, как ядро C будет выполнено.
Например, если ядро C в файле CU имеет следующую подпись:
void reallySimple( float * pInOut, float c )
соответствующий объект ядра (k
) в MATLAB имеет следующие свойства:
MaxNumLHSArguments: 1 NumRHSArguments: 2 ArgumentTypes: {'inout single vector' 'in single scalar'}
Поэтому использовать объект ядра из этого кода с feval
, необходимо обеспечить feval
два входных параметра (в дополнение к объекту ядра), и можно использовать один выходной аргумент:
y = feval(k,x1,x2)
Входные значения x1
и x2
соответствуйте pInOut
и c
в прототипе функции C. Выходной аргумент y
соответствует значению pInOut
в прототипе функции C после того, как выполнилось ядро C.
Следующее является немного более сложным примером, который показывает комбинацию указателей неconst и const:
void moreComplicated( const float * pIn, float * pInOut1, float * pInOut2 )
Соответствующий объект ядра в MATLAB затем имеет свойства:
MaxNumLHSArguments: 2 NumRHSArguments: 3 ArgumentTypes: {'in single vector' 'inout single vector' 'inout single vector'}
Можно использовать feval
на ядре этого кода (k
) с синтаксисом:
[y1,y2] = feval(k,x1,x2,x3)
Эти три входных параметра x1
x2
, и x3
, соответствуйте трем аргументам, которые передаются в функцию C. Выходные аргументы y1
и y2
, соответствуйте значениям pInOut1
и pInOut2
после того, как ядро C выполнилось.
Этот пример добавляет два, удваивается вместе в графическом процессоре. Вы должны иметь Инструментарий CUDA NVIDIA, установленный, и иметь CUDA-способные драйверы для вашего устройства.
Код CU, чтобы сделать это следующие.
__global__ void add1( double * pi, double c ) { *pi += c; }
Направляющий __global__
указывает, что это - точка входа к ядру. Код использует указатель, чтобы отослать результат в pi
, который является и входом и выходом. Поместите этот код в файл под названием test.cu
в текущем каталоге.
Скомпилируйте код CU в командной строке интерпретатора, чтобы сгенерировать файл PTX под названием test.ptx
.
nvcc -ptx test.cu
Создайте ядро в MATLAB. В настоящее время этот файл PTX только имеет одну запись, таким образом, вы не должны задавать его. Если бы необходимо было вставить больше ядер, вы задали бы add1
как запись.
k = parallel.gpu.CUDAKernel('test.ptx','test.cu');
Запустите ядро с двумя числовыми входными параметрами. По умолчанию ядро работает на одном потоке.
result = feval(k,2,3)
result = 5
Этот пример расширяет предыдущий, чтобы добавить два вектора вместе. Для простоты примите, что существует точно то же количество потоков как элементы в векторах и что существует только один блок потока.
Код CU немного отличается из последнего примера. Оба входных параметров являются указателями, и каждый является постоянным, потому что вы не изменяете его. Каждый поток просто добавит элементы в своем списке веток. Список веток должен удаться, какой элемент этот поток должен добавить. (Получение их распараллеливает - и специфичные для блока значения очень общий шаблон в программировании CUDA.)
__global__ void add2( double * v1, const double * v2 ) { int idx = threadIdx.x; v1[idx] += v2[idx]; }
Сохраните этот код в файле test.cu
.
Скомпилируйте как перед использованием nvcc
.
nvcc -ptx test.cu
Если бы этот код был помещен в тот же файл CU наряду с кодом первого примера, необходимо задать имя точки входа на этот раз, чтобы отличить его.
k = parallel.gpu.CUDAKernel('test.ptx','test.cu','add2');
Прежде чем вы запустите ядро, определите номер потоков правильно для векторов, вы хотите добавить.
N = 128; k.ThreadBlockSize = N; in1 = ones(N,1,'gpuArray'); in2 = ones(N,1,'gpuArray'); result = feval(k,in1,in2);
Для примера, который показывает, как работать с CUDA, и обеспечивает CU и файлы PTX для вас, чтобы экспериментировать с, смотрите Иллюстрирование Трех Подходов к Вычислению графического процессора: Множество Мандельброта.