CUDA всем хороша, пока под рукой есть видеокарта от Nvidia. Но что делать, когда на любимом ноутбуке нет Nvidia видеокарты? Или нужно вести разработку в виртуальной машине?

Я постараюсь рассмотреть в этой статье такое решение, как фреймворк rCUDA (Remote CUDA), который поможет, когда Nvidia видеокарта есть, но установлена не в той машине, на которой предполагается запуск CUDA приложений. Тем, кому это интересно, добро пожаловать под кат.

Что такое rCUDA

rCUDA (Remote CUDA) — фреймворк, реализующий CUDA API, позволяющий использовать для CUDA вычислений видеокарту, расположенную на удалённой машине, не внося никаких изменений в ваш код. Разработан в политехническом университете Валенсии (rcuda-team).

Ограничения

На данный момент поддерживаются только GNU/Linux системы, однако разработчики обещают поддержку Windows в будущем. Текущая версия rCUDA, 18.03beta, совместима с CUDA 5-8, то есть CUDA 9 не поддерживается. Разработчиками заявлена полная совместимость с CUDA API, за исключением графики.

Возможные сценарии использования

1. Запуск CUDA приложений в виртуальной машине тогда, когда проброс видеокарты неудобен или невозможен, например, когда видеокарта занята хостом, или когда виртуальных машин больше одной.
2. Ноутбук без дискретной видеокарты.
3. Желание использовать несколько видеокарт (кластеризация). Теоретически, можно использовать все имеющиеся в команде видеокарты, в том числе совместо.

Краткая инструкция

Тестовая конфигурация

Тестирование проводилось на следующей конфигурации:

Сервер:
Ubuntu 16.04, GeForce GTX 660

Клиент:
Виртуальная машина с Ubuntu 16.04 на ноутбуке без дискретной видеокарты.

Получение rCUDA

Cамый сложный этап. К сожалению, на данный момент единственный способ получить свой экземпляр этого фреймворка — заполнить соответствующую форму запроса на официальном сайте. Впрочем, разработчики обещают отвечать в течение 1-2 дней. В моём случае мне прислали дистрибутив в тот же день.

Установка CUDA

Для начала необходимо установить CUDA Toolkit на сервере и клиенте (даже если на клиенте нет nvidia видеокарты). Для этого можно скачать его с официального сайта или использовать репозиторий. Главное, использовать версию не выше 8. В данном примере используется установщик .run с оффициального сайта.

chmod +x cuda_8.0.61_375.26_linux.run
./cuda_8.0.61_375.26_linux.run

Важно! На клиенте следует отказаться от установки nvidia драйвера. По умолчанию CUDA Toolkit будет доступен по адресу /usr/local/cuda/. Установите CUDA Samples, они понадобятся.

Установка rCUDA

Распакуем полученный от разработчиков архив в нашу домашнюю директорию на сервере и на клиенте.

tar -xvf rCUDA*.tgz -C ~/
mv ~/rCUDA* ~/rCUDA

Проделать эти действия нужно как на сервере, так и на клиенте.

Запуск демона rCUDA на сервере

export PATH=$PATH/usr/local/cuda/bin
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib64:/home/<XXX>/rCUDA/lib/cudnn
cd ~/rCUDA/bin
./rCUDAd  

Замените < XXX> на имя вашего пользователя. Используйте ./rCUDAd -iv, если хотите видеть подробный вывод.

Настройка клиента

Откроем на клиенте терминал, в котором в дальнейшем будем запускать CUDA код. На стороне клиента нам необходимо "подменить" стандартные библиотеки CUDA на библиотеки rCUDA, для чего добавим соответствующие пути в переменную среды LD_LIBRARY_PATH. Также нам необходимо указать количество серверов и их адреса (в моём примере он будет один).

export PATH=$PATH/usr/local/cuda/bin  
export LD_LIBRARY_PATH=/home/<XXX>/rCUDA/lib/:$LD_LIBRARY_PATH 
export RCUDA_DEVICE_COUNT=1  # укажем количество видеокарт (серверов), их может быть несколько 
export RCUDA_DEVICE_0=<IP АДРЕС СЕРВЕРА>:0 # укажем адрес первого сервера

Сборка и запуск

Попробуем собрать и запустить несколько примеров.

Пример 1

Начнём с простого, с deviceQuery — примера, который просто выведет нам параметры CUDA совместимого устройства, то есть в нашем случае удалённого GTX660.

cd <YYY>/NVIDIA_CUDA-8.0_Samples/1_Utilities/deviceQuery
make EXTRA_NVCCFLAGS=--cudart=shared

Важно! Без EXTRA_NVCCFLAGS=--cudart=shared чуда не получится
Замените <YYY> на путь, который вы указали для CUDA Samples при установке CUDA.

Запустим собранный пример:

./deviceQuery

Если вы всё сделали правильно, результат будет примерно таким:

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 660"
  CUDA Driver Version / Runtime Version          9.0 / 8.0
  CUDA Capability Major/Minor version number:    3.0
  Total amount of global memory:                 1994 MBytes (2090991616 bytes)
  ( 5) Multiprocessors, (192) CUDA Cores/MP:     960 CUDA Cores
  GPU Max Clock rate:                            1072 MHz (1.07 GHz)
  Memory Clock rate:                             3004 Mhz
  Memory Bus Width:                              192-bit
  L2 Cache Size:                                 393216 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 9.0, CUDA Runtime Version = 8.0, NumDevs = 1, Device0 = GeForce GTX 660
Result = PASS

Самое главное, что мы должны увидеть:

Device0 = GeForce GTX 660
Result = PASS

Отлично! Нам удалось собрать и запустить CUDA приложение на машине без дискретной видеокарты, использовав для этого видеокарту, установленную на удалённом сервере.

Важно! Если вывод приложения начинается со строк вида:

mlock error: Cannot allocate memory
rCUDA warning: 1007.461
mlock error: Cannot allocate memory 

значит необходимо добавить на сервере и на клиенте в файл "/etc/security/limits.conf" следующие строки:

*            hard   memlock           unlimited
*            soft    memlock           unlimited 

Таким образом, вы разрешите всем пользователям (*) неограниченное (unlimited) блокирование памяти (memlock). Еще лучше будет заменить * на нужного пользователя, а вместо unlimited подобрать менее жирные права.

Пример 2

Теперь попробуем что-то поинтереснее. Протестируем реализацию скалярного произведения векторов с использованием разделяемой памяти и синхронизации ("Технология CUDA в примерах" Сандерс Дж. Кэндрот Э. 5.3.1).

В данном примере мы рассчитаем скалярное произведение двух векторов размерностью 33 * 1024, сравнивая ответ с результатом, полученным на CPU.

#include <stdio.h>

#define imin(a,b) (a<b?a:b)

const int N = 33 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid = imin(32, (N+threadsPerBlock-1) / threadsPerBlock);

__global__ void dot(float* a, float* b, float* c) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float temp = 0;
    while (tid < N){
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }

    // set the cache values
    cache[cacheIndex] = temp;

    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0){
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}

int main (void) {
    float *a, *b, c, *partial_c;
    float *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the cpu side
    a = (float*)malloc(N*sizeof(float));
    b = (float*)malloc(N*sizeof(float));
    partial_c = (float*)malloc(blocksPerGrid*sizeof(float));

    // allocate the memory on the gpu
    cudaMalloc((void**)&dev_a, N*sizeof(float));
    cudaMalloc((void**)&dev_b, N*sizeof(float));
    cudaMalloc((void**)&dev_partial_c, blocksPerGrid*sizeof(float));

    // fill in the host memory with data
    for(int i=0; i<N; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    // copy the arrays 'a' and 'b' to the gpu
    cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice);

    dot<<<blocksPerGrid, threadsPerBlock>>>(dev_a, dev_b, dev_partial_c);

    // copy the array 'c' back from the gpu to the cpu
    cudaMemcpy(partial_c,dev_partial_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost);

    // finish up on the cpu side
    c = 0;
    for(int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    #define sum_squares(x) (x*(x+1)*(2*x+1)/6)
    printf("GPU - %.6g \nCPU - %.6g\n", c, 2*sum_squares((float)(N-1)));

    // free memory on the gpu side
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_partial_c);

    // free memory on the cpu side
    free(a);
    free(b);
    free(partial_c);
}

Сборка и запуск:

/usr/local/cuda/bin/nvcc  --cudart=shared dotProd.cu -o dotProd
./dotProd

Такой результат говорит нам, что всё у нас хорошо:

GPU — 2.57236e+13
CPU — 2.57236e+13

Пример 3

Запустим еще один стандартный тест CUDA- matrixMulCUBLAS (перемножение матриц).

cd < YYY>/NVIDIA_CUDA-8.0_Samples/0_Simple/matrixMulCUBLAS
make EXTRA_NVCCFLAGS=--cudart=shared
./matrixMulCUBLAS

Интересное нам:

Performance= 436.24 GFlop/s,
Comparing CUBLAS Matrix Multiply with CPU results: PASS

Безопасность

Я не нашёл в документации к rCUDA упоминания о каком-либо способе авторизации. Думаю, на данный момент самое простое, что можно сделать, это открыть доступ к нужному порту (8308) только с определённого адреса.

При помощи iptables это будет выглядеть так:

iptables -A INPUT -m state --state NEW -p tcp -s <адрес клиента> --dport 8308 -j ACCEPT

В остальном оставляю вопрос безопасности за рамками данного поста.