CONFLUX: GPGPU ДЛЯ.NET Евгений Бурмако Андрей Воронович. - презентация

1 CONFLUX: GPGPU ДЛЯ.NET Евгений Бурмако Андрей Воронович

2 Видеокарты сегодня На борту – десятки/сотни ALU на частоте более 1 GHz. В пике – 1 TFLOPS (и >100 GFLOPS двойной точности) API – произвольный доступ к памяти, структуры данных, указатели, подпрограммы. Возраст API – почти четыре года, несколько поколений графических процессоров.

3 С точки зрения программиста Современные модели программирования GPU (CUDA, AMD Stream, OpenCL, DirectCompute): Параллельный алгоритм задается парой: 1) ядро (итерация цикла), 2) границы итерации. Ядро компилируется драйвером. На основе границ итерации создается решетка вычислительных потоков. Входные данные копируются в видеопамять. Запускается выполнение задачи. Результат копируется в оперативную память.

4 Пример: SAXPY на CUDA __global__ void Saxpy(float a, float* X, float* Y) { int i = blockDim.x * blockIdx.x + threadIdx.x; Y[i] = a * X[i] + Y[i]; } cudaMemcpy(X, hX, cudaMemcpyHostToDevice); cudaMemcpy(Y, hY, cudaMemcpyHostToDevice); Saxpy >>(a, hX, hY); cudaMemcpy(hY, Y, cudaMemcpyDeviceToHost);

5 Вопрос к знатокам

6 Отвечает Александр Друзь

7 По факту Brahma: Структуры данных: data parallel array. Вычисления: выражения C#, LINQ-комбинаторы. Accelerator v2: Структуры данных: data parallel array. Вычисления: арифметические операторы, набор предопределенных функций. Это работает для многих алгоритмов. Но что, если есть ветвления или нерегулярный доступ к памяти?

8 А вот что saxpy void Saxpy(float a, float* X, float* Y) { int i = blockDim.x * blockIdx.x + threadIdx.x; Y[i] = a * X[i] + Y[i]; }; nvcuda.cuModuleLoadDataEx(saxpy); nvcuda.cuMemcpyHtoD(X, Y); nvcuda.cuParamSeti(a, X, Y); nvcuda.cuLaunchGrid(256, (N + 255) / 256); nvcuda.cuMemcpyDtoH(Y);

9 Конфлакс Ядра пишутся на С#: поддерживаются структуры данных, локальные переменные, ветвления, циклы. float a; float[] x; [Result] float[] y; var i = GlobalIdx.X; y[i] = a * x[i] + y[i];

10 Конфлакс Не требует явного общения с неуправляемым кодом, позволяет работать с родными типами данных.NET. float[] x, y; var cfg = new CudaConfig(); var kernel = cfg.Configure (); y = kernel.Execute(a, x, y);

11 Как это работает? Front end: декомпилятор C#. Преобразование AST: инлайн вызываемых методов, деструктуризация классов и массивов, отображение вычислительных операций. Back end: генераторы PTX и многоядерного IL. Привязка к драйверу nvcuda, который умеет исполнять программы на ассемблере.

12 Успехи Альфа-версия. Умеет вычислять hello-world параллельных вычислений: умножение матриц. За вычетом [на текущий момент] высоких издержек на JIT-компиляцию идея оправдывает себя даже для наивного кодогенератора: 1x CPU < 2x CPU

13 Демонстрация

14 Следующие шаги Оптимизации для графических процессоров (лесенка для оптимальной пропускной способности при транспозиции матриц). Полиэдральная модель оптимизации циклов (конфигурируется относительно иерархии и размеров кэшей, есть линейные эвристики, оптимизирующие локальность данных в вычислительной решетке). Исполнение на кластере (следующий шаг после полиэдральной модели: добавляется виртуальный уровень кэша – вычислительный узел).

15 Заключение Ресурсы: coming soon

16 СПАСИБО!