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

Скомпилируйте файл PTX из файла CU

Если у вас есть файл CU, вы хотите выполниться на графическом процессоре, необходимо сначала скомпилировать его, чтобы создать файл PTX. Один способ сделать это с компилятором nvcc в NVIDIA^® CUDA^® Toolkit. Например, если ваш файл CU называется myfun.cu, можно создать скомпилированный файл PTX с командой интерпретатора:

nvcc -ptx myfun.cu

Это генерирует файл с именем myfun.ptx.

Создайте объект CUDAKernel с входом файла CU

С файлом .cu и файлом .ptx можно создать объект CUDAKernel в MATLAB, который можно затем использовать, чтобы оценить ядро:

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');

Создайте объект CUDAKernel с прототипным входом C

Если у вас нет файла 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 ...)

При переводе определения ядра в C в MATLAB:

Эти правила имеют некоторые последствия. Самое известное - то, что каждый вывод от ядра должен обязательно также быть входом к ядру, поскольку вход позволяет пользователю задавать размер вывода (который следует из неспособности выделить память на графическом процессоре).

Свойства объектов CUDAKernel

Когда вы создаете объект ядра без останавливающейся точки с запятой, или когда вы вводите переменную объекта в командной строке, 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, чтобы сгенерировать ядро.

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

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 путем установки двух из его свойств объектов:

Значением по умолчанию для обоих из этих свойств является [1 1 1], но предположите, что вы хотите использовать 500 потоков, чтобы запустить поэлементные операции на векторах 500 элементов параллельно. Простой способ достигнуть этого состоит в том, чтобы создать ваш CUDAKernel и установить его свойства соответственно:

k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu');
k.ThreadBlockSize = [500,1,1];

Обычно вы устанавливаете сетку и распараллеливаете размеры блока на основе размеров ваших входных параметров. Для получения информации об иерархии потока, и сетках нескольких-размерностей и блоках, смотрите NVIDIA CUDA C Руководство по программированию.

Используйте переменные рабочей области

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

k = parallel.gpu.CUDAKernel('conv.ptx','conv.cu');

result = feval(k,rand(100,1),rand(100,1));

Даже если входные параметры являются константами или переменными для данных о рабочем пространстве MATLAB, выводом является gpuArray.

Используйте 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 и файлами PTX

Для примера, который показывает, как работать с CUDA, и обеспечивает CU и файлы PTX для вас, чтобы экспериментировать с, смотрите Иллюстрирование Трех Подходов к Вычислению графического процессора: Множество Мандельброта.