[Перевод - recovery mode ] Программирование GPU на Java02.05.2020 20:48

Получение доступ к GPU из Java раскрывает огромную мощь. Здесь рассказывается как GPU работает и как получить доступ из Java.

Программирование устройства графического процессора (GPU) является заоблачным миром для Java программистов. Это понятно, так как обычные задачи для Java не подходят для GPU. Тем не менее, GPU обладают терафлопсами производительности, так давайте исследуем их возможности.
Для того чтобы сделать топик доступным, я потрачу некоторое время объясняя архитектуру GPU вместе с небольшой историей, которая облегчит погружение в программирование железа.

Однажды мне показали отличия GPU от CPU вычислений, я покажу как использовать GPU в мире Java. Наконец, я опишу главные фреймворки и библиотеки доступные для написания кода на Java и запуска их на GPU, и я приведу некоторые примеры кода.

Немного бекграунда

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

Суть 2D/3D обработки состоит в основном в манипулировании матрицами, это может быть управляемо с помощью распределенного подхода. Что будет эффективным подходом для обработки изображений? Ответим на это, давайте сравним стандартную архитектуру CPU (показанную на рис. 1.) и GPU.

Рис. 1. Блоки архитектуры CPU

В CPU, фактические элементы обработки — регистры, арифметико-логический блок (ALU) и контексты выполнения — лишь маленькие части всей системы. Для ускорения нерегулярных расчетов, поступающих в непредсказуемом порядке, здесь присутствует большой, быстрый, и дорогой кеш; различные виды сборщиков; и предсказатели ветвлений.

Вам не нужно все это на GPU, потому что получены данные в предсказуемой манере, и GPU производит очень ограниченное множество операций над данными. Таким образом, возможно сделать их очень маленькими и не дорогой процессор с блочной архитектурой, похожую на эту, показан на рис. 2.

Рис. 2. Блоковая архитектура для простого ядра GPU

Потому что такие процессоры дешевле и обрабатываемые данные в них в параллельных кусках, это просто заставить многие из них работать параллельно. Это разработано, ссылаясь на множественные инструкции, множественные данные или MIMD (произносится «mim-dee»).

Второй подход базируется на факте, что часто одиночная инструкция, применяется к множественным частям данных. Это известно как единичная инструкция, множественные данные или SIMD (произносится «sim-dee»). В этом дизайне, единственный GPU содержит множественные ALU и контексты выполнения, маленькие области, переданные на разделяемые данные контекста, как показано на рис. 3.

Рис. 3. Сравнение MIMD-стиля архитектуры GPU блоков (с лево), с SIMD дизайном (с право)

Смешивая SIMD и MIMD обработку обеспечивается максимальная пропускная способность, которую я обойду стороной. В этом дизайне, вы имеете множественные SIMD процессоры, работающие параллельно, как на рис. 4.

Рис. 4. Работающие множественные SIMD процессоры в параллели; здесь 16 ядер с 128 ALU

Так как вы имеете связку маленьких, простых процессоров, вы можете программировать их для получения специального эффекта в выводе.

Запуская программы на GPU

Большинство ранних графических эффектов в играх было действительно жестко запрограммированных маленьких программ выполняющихся на GPU и применялись к потокам данных из CPU.

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

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

Каждый продавец GPU имел свой собственный язык программирования и инфраструктуру для создания shaders для их архитектуры. На этом подходе было создано множество платформ.

Главные из них:

• DirectCompute: частные shader язык/API от Microsoft, которые являются частью Direct3D, начиная с DirectX 10.
• AMD FireStream: частные технологии ATI/Radeon, которые устарели по мнению AMD.
• OpenACC: мультивендор-консорциум, решение для параллельных вычислений
• С++ AMP: частная библиотека Microsoft для параллельных данных на C++
• CUDA: частная платформа Nvidia, которая использует подмножество языка С
• OpenСL: общий стандарт, оригинально разработанный Apple, но сейчас не управляем консорциумом Khronos Group

Большую часть времени, работа с GPU это низко уровневое программирование. Для того чтобы сделать это немного более понятным для разработчиков, для кодирования, были обеспеченны несколько абстракций. Самая известная это DirectX, от Microsoft, и OpenGL, от Khronos Group. Эти API для написания высокоуровневого кода, который затем может быть затем упрощен для GPU, более семантически, для программиста.

Насколько я знаю, не существует инфраструктуры Java для DirectX, но есть хорошее решение для OpenGL. JSR 231 запустилось в 2002 и адресовано GPU программистам, но оно было заброшено в 2008 и поддерживает только OpenGL 2.0.

Поддержка OpenGL продолжается в независимом проекте JOCL, (который также поддерживает OpenCL), и он доступен аудитории. Таким образом, знаменитая игра Minecraft была написана с использованием JOCL.

Приход GPGPU

До сих пор, Java и GPU не имели точек соприкосновения, хотя они должны быть. Java часто используется на предприятиях, в науке о данных, и в финансовом секторе, где много вычислений и где нужно много вычислительных сил. Это о том, как идея general-purpose GPU (GPGPU). Идея использование GPU по этому пути начали, когда производители видео адаптеров начали давать доступ к программному буферу кадров, давая разработчикам считывать содержимое. Некоторые хакеры определили, что они могут использовать всю силу GPU для универсальных вычислений.
Рецепт был таким:

1. Закодировать данные как растровый массив.
2. Написать shaders для обработки их.
3. Отправьте их обоих на видеокарту.
4. Получить результат из буфера кадра
5. Декодировать данные из растрового массива.

Это очень простое объяснение. Я не уверен будет ли это работать в продакшене, но это действительно работает.

Затем начались множественные исследования от Стенфордского Института для упрощения использования GPU. В 2005 они сделали BrookGPU, которая была маленькой экосистемой, которая включала язык программирования, компилятор, и среду запуска.

BrookGPU компилировала программы написанные на языке программирования потоков Brook, который был вариантом ANSI C. Он может быть нацелен на OpenGL v1.3 +, DirectX v9 + или AMD Close to Metal для вычислительной серверной части, и он запускается на Microsoft Windows и Linux. Для отладки, BrookGPU может также симулировать виртуальную графическую карту на CPU.
Однако это не взлетело, из-за оборудования, доступного в то время. В мире GPGPU, тебе надо копировать данные на устройство (в это контексте, устройство ссылается на GPU и устройстве, на котором он расположен), ждать GPU для вычисления данных, и затем копировать данные обратно в управляющую программу. Это создает множество задержек. И в середине 2000-х, когда проект был в активной разработке, эти задержки также исключали интенсивное использование GPU для базовых вычислений.

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

Теперь, когда я вам все рассказал, давайте проверим две самые удачные технологии для вычисления на GPU — OpenCL и CUDA — смотри также как Java работает с ними.

OpenCL и Java

Как и другие инфраструктурные пакеты, OpenCL обеспечивает базовую имплементацию на C. Это технически доступно с помощью Java Native Interface (JNI) или Java Native Access (JNA), но такой подход будет слишком тяжёлым для большинства разработчиков.

К счастью, эта работа была уже сделана несколькими библиотеками: JOCL, JogAmp, и JavaCL. К несчастью, JavaCL стал мертвым проектом. Но проект JOCL жив и очень даже адаптирован. Я буду использовать его для следующих примеров.

Но сначала я должен объяснить, что из себя представляет OpenCL. Я упоминал ранее, что OpenCL обеспечивает очень базовую модель, подходящую для программирования всех видов устройств — не только GPU и CPU, но даже DSP процессоры и FPGA.

Давайте посмотрим на самый простой пример: складывание векторов, наверное, самый яркий и простой пример. Вы имеете два массива чисел для сложение и один для результата. Вы берете элемент из первого массива и элемент из второго массива, и затем вы кладете сумму в массив результатов, показано на рис. 5.

Рис. 5. Сложение элементов двух массивов и сохранение суммы в результирующий массив

Как вы можете видеть, операция очень согласованная и тем не менее распределяемая. Вы можете запихнуть каждую операцию сложения в разные GPU ядра. Это значит, что, если вы имеете 2048 ядер, как на Nvidia 1080, вы можете выполнить 2048 операций сложения одновременно. Это означает что здесь потенциальные терафлопсы компьютерной мощи ждут вас. Этот код для массива из 10 миллионов чисел взят с сайта JOCL:

public class ArrayGPU {
    /**
     * The source code of the OpenCL program 
     */
    private static String programSource =
        "__kernel void "+
        "sampleKernel(__global const float *a,"+
        "             __global const float *b,"+
        "             __global float *c)"+
        "{"+
        "    int gid = get_global_id(0);"+
        "    c[gid] = a[gid] + b[gid];"+
        "}";
    
    public static void main(String args[])
    {
        int n = 10_000_000;
        float srcArrayA[] = new float[n];
        float srcArrayB[] = new float[n];
        float dstArray[] = new float[n];
        for (int i=0; i

__kernel void MatrixMul_kernel_basic(int dim,
                  __global float *A,
                  __global float *B,
                  __global float *C){

    int iCol = get_global_id(0);
    int iRow = get_global_id(1);
    float result = 0.0;
    for(int i=0; i< dim; ++i)
    {
          result +=
          A[iRow*dim + i]*B[i*dim + iCol];
    }
    C[iRow*dim + iCol] = result;
}

#define VECTOR_SIZE 4    
__kernel void MatrixMul_kernel_basic_vector4(
    size_t dim, // dimension is in single floats
    const float4 *A,
    const float4 *B,
    float4 *C)
{
    size_t globalIdx = get_global_id(0);
    size_t globalIdy = get_global_id(1);
    float4 resultVec = (float4){ 0, 0, 0, 0 };
    size_t dimVec = dim / 4;
    for(size_t i = 0; i < dimVec; ++i) {
        float4 Avector = A[dimVec * globalIdy + i];
        float4 Bvector[4];
        Bvector[0] = B[dimVec * (i * 4 + 0) + globalIdx];
        Bvector[1] = B[dimVec * (i * 4 + 1) + globalIdx];
        Bvector[2] = B[dimVec * (i * 4 + 2) + globalIdx];
        Bvector[3] = B[dimVec * (i * 4 + 3) + globalIdx];
        resultVec += Avector[0] * Bvector[0];
        resultVec += Avector[1] * Bvector[1];
        resultVec += Avector[2] * Bvector[2];
        resultVec += Avector[3] * Bvector[3];
    }

    C[dimVec * globalIdy + globalIdx] = resultVec;
}

...
int n = 100;
curandGenerator generator = new curandGenerator();
float hostData[] = new float[n];
Pointer deviceData = new Pointer();
cudaMalloc(deviceData, n * Sizeof.FLOAT);
curandCreateGenerator(generator, CURAND_RNG_PSEUDO_DEFAULT); 
curandSetPseudoRandomGeneratorSeed(generator, 1234);
curandGenerateUniform(generator, deviceData, n);
cudaMemcpy(Pointer.to(hostData), deviceData, 
        n * Sizeof.FLOAT, cudaMemcpyDeviceToHost);
System.out.println(Arrays.toString(hostData));
curandDestroyGenerator(generator);
cudaFree(deviceData);
...

public static void main(String[] _args) {
    final int size = 512;
    final float[] a = new float[size];
    final float[] b = new float[size];

    /* fill the arrays with random values */
    for (int i = 0; i < size; i++){
        a[i] = (float) (Math.random() * 100);
        b[i] = (float) (Math.random() * 100);
    }
    final float[] sum = new float[size];

    Kernel kernel = new Kernel(){
        @Override public void run() {
I           int gid = getGlobalId();
            sum[gid] = a[gid] + b[gid];
        }
    };

    kernel.execute(Range.create(size));
    for(int i = 0; i < size; i++) {
        System.out.printf("%6.2f + %6.2f = %8.2f\n", a[i], b[i], sum[i])
    }
    kernel.dispose();
}