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

Компиляция файла PTX из файла CU

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

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:

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

Свойства объекта 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 для версии float, и 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 для выполнения на GPU. У вас есть ядро, которое выполняет свертку на двух векторах; загрузить и запустить его с двумя случайными входными векторами:

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 и non-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.

Добавить два числа

В этом примере в GPU добавляются два двойника. Необходимо установить набор инструментов NVIDIA CUDA Toolkit и иметь драйверы, поддерживающие CUDA.

Добавить два вектора

Этот пример расширяет предыдущий, чтобы добавить два вектора вместе. Для простоты предположим, что существует точно такое же количество нитей, как и элементов в векторах, и что существует только один блок нитей.

Пример с файлами CU и PTX

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