CUDA e GPU remota

O CUDA é bom para todos, desde que haja uma placa de vídeo da Nvidia em mãos. Mas o que fazer quando não há placa de vídeo Nvidia no seu laptop favorito? Ou você precisa realizar o desenvolvimento em uma máquina virtual?

Tentarei considerar neste artigo uma solução como a estrutura rCUDA (Remote CUDA), que ajudará quando houver uma placa de vídeo Nvidia, mas ela não está instalada na máquina em que os aplicativos CUDA devem ser lançados. Para aqueles que estão interessados, bem-vindo ao gato.

O que é rCUDA

O rCUDA (CUDA remoto) é uma estrutura que implementa a API CUDA, permitindo que você use uma placa de vídeo localizada na máquina remota para a computação CUDA sem fazer alterações no seu código. Desenvolvido na Universidade Politécnica de Valência ( rcuda-team ).

Limitações

Atualmente, apenas os sistemas GNU / Linux são suportados, no entanto, os desenvolvedores prometem suporte ao Windows no futuro. A versão atual do rCUDA, 18.03beta, é compatível com o CUDA 5-8, ou seja, o CUDA 9 não é suportado. Os desenvolvedores declararam total compatibilidade com a API CUDA, com exceção dos gráficos.

Possíveis casos de uso

1. A execução de aplicativos CUDA em uma máquina virtual ao encaminhar uma placa de vídeo é inconveniente ou impossível, por exemplo, quando a placa de vídeo está ocupada por um host ou quando há mais de uma máquina virtual.
2. Laptop sem placa gráfica discreta.
3. O desejo de usar várias placas de vídeo (clustering). Teoricamente, você pode usar todas as placas de vídeo disponíveis na equipe, inclusive em conjunto.

Breve instrução

Configuração de teste

O teste foi realizado na seguinte configuração:

Servidor:
Ubuntu 16.04, GeForce GTX 660

Cliente:
Uma máquina virtual com Ubuntu 16.04 em um laptop sem uma placa gráfica discreta.

Obtendo o rCUDA

A fase mais difícil. Infelizmente, no momento, a única maneira de obter sua cópia dessa estrutura é preencher o formulário de solicitação apropriado no site oficial. No entanto, os desenvolvedores prometem responder dentro de 1-2 dias. No meu caso, eles me enviaram uma distribuição no mesmo dia.

Instalar CUDA

Primeiro, você precisa instalar o CUDA Toolkit no servidor e no cliente (mesmo se o cliente não tiver uma placa de vídeo nvidia). Para fazer isso, você pode baixá-lo do site oficial ou usar o repositório. O principal é usar a versão não superior a 8. Neste exemplo, o instalador .run do site oficial é usado .

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

Importante! No cliente, você deve recusar a instalação do driver da nvidia. Por padrão, o CUDA Toolkit estará disponível em / usr / local / cuda /. Instale as amostras CUDA, você precisará delas.

Instale o rCUDA

Descompactaremos o arquivo recebido dos desenvolvedores em nosso diretório pessoal no servidor e no cliente.

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

Você precisa executar essas ações no servidor e no cliente.

Iniciando o daemon rCUDA no servidor

 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 

Substitua <XXX> pelo seu nome de usuário. Use ./rCUDAd -iv se desejar ver uma saída detalhada.

Configuração do cliente

Vamos abrir o terminal no cliente, no qual executaremos o código CUDA no futuro. No lado do cliente, precisamos "substituir" as bibliotecas CUDA padrão pelas bibliotecas rCUDA, para as quais adicionamos os caminhos apropriados à variável de ambiente LD_LIBRARY_PATH. Também precisamos especificar o número de servidores e seus endereços (no meu exemplo, será um).

 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 #     

Montagem e lançamento

Vamos tentar criar e executar alguns exemplos.

Exemplo 1

Vamos começar com um exemplo simples de deviceQuery que simplesmente exibe as configurações CUDA para um dispositivo compatível, ou seja, no nosso caso, o GTX660 remoto.

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

Importante! Sem EXTRA_NVCCFLAGS = - cudart = compartilhado, o milagre não funcionará
Substitua <YYY> pelo caminho especificado para amostras CUDA ao instalar o CUDA.

Execute o exemplo montado:

 ./deviceQuery 

Se você fez tudo corretamente, o resultado será algo como isto:

 ./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 

A coisa mais importante que devemos ver:

Device0 = GeForce GTX 660
Resultado = PASS

Ótimo! Conseguimos criar e executar o aplicativo CUDA em uma máquina sem uma placa gráfica discreta, usando para isso uma placa de vídeo instalada em um servidor remoto.

Importante! Se a saída do aplicativo começar com linhas do formulário:

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

significa que é necessário adicionar as seguintes linhas ao arquivo "/etc/security/limits.conf" no servidor e no cliente:

 * hard memlock unlimited * soft memlock unlimited 

Assim, você permitirá a todos os usuários (*) memória de bloqueio ilimitada (ilimitada) (memlock). Seria ainda melhor substituir * pelo usuário desejado e, em vez de ilimitado, escolher direitos com menos gordura.

Exemplo 2

Agora vamos tentar algo mais interessante. Testaremos a implementação do produto escalar de vetores usando memória compartilhada e sincronização ("CUDA Technology in Examples" Sanders J. Kendrot E. 5.3.1).

Neste exemplo, calculamos o produto escalar de dois vetores da dimensão 33 * 1024, comparando a resposta com o resultado obtido na 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); } 

Crie e execute:

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

Esse resultado nos diz que está tudo bem conosco:

GPU - 2.57236e + 13
CPU - 2.57236e + 13

Exemplo 3

Execute outro teste padrão CUDA-matrixMulCUBLAS (multiplicação de matriz).

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

Interessante para nós:

Desempenho = 436,24 GFlop / s,
Comparando os resultados da matriz CUBLAS Multiply com CPU: PASS

Segurança

Não encontrei menção de nenhum método de autorização na documentação do rCUDA. Acho que, no momento, a coisa mais simples que pode ser feita é abrir o acesso à porta desejada (8308) apenas a partir de um endereço específico.

Usando o iptables, ficará assim:

 iptables -A INPUT -m state --state NEW -p tcp -s < > --dport 8308 -j ACCEPT 

De resto, deixo a questão de segurança além do escopo deste post.