(c) Cleuton Sampaio 2019
A programação CUDA utilizando Go é um pouco mais complexa do que em outras linguagens. Apesar de existirem alguns pacotes excelentes, como o mumax, a documentação é ruim, carece de exemplos e são difíceis de utilizar.
CUDA é para C, então, a melhor alternativa é usar o Command cgo e invocar uma função externa com o seu Kernel cuda. É o que farei nesse exemplo, onde multiplico duas matrizes usando CUDA.
Se você quiser saber mais sobre programação CUDA, leia o meu artigo.
Criei um Kernel simples que possui a função Kernel e uma helper function para ser chamada externamente. Note que eu usei extern C pois é assim que o cgo invoca funções:
#include <stdio.h>
#include <cuda.h>
__global__ void vecmul(float *A, float* B, float *C, int size)
{
// Row and Column indexes:
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
// Are they bellow the maximum?
if (col < size && row < size) {
float result = 0;
for(int ix=0;ix<size;ix++) {
result += A[row*size+ix]*B[ix*size+col];
}
C[row*size+col] = result;
}
}
extern "C" {
void maxmul(float *A, float* B, float *C, int size) {
int total = size*size;
// Allocate device memory:
float* gpu_A;
float* gpu_B;
float* gpu_C;
int msize = total * sizeof(float);
cudaMalloc((void**)&gpu_A, msize);
cudaMemcpy(gpu_A,A,msize,cudaMemcpyHostToDevice);
cudaMalloc((void**)&gpu_B, msize);
cudaMemcpy(gpu_B,B,msize,cudaMemcpyHostToDevice);
cudaMalloc((void**)&gpu_C,msize);
// Blocks & grids:
dim3 blocks(size,size);
dim3 grid(1,1);
// Call the kernel:
vecmul<<<grid,blocks>>>(gpu_A,gpu_B,gpu_C,size);
// Get the result Matrix:
cudaMemcpy(C,gpu_C,msize,cudaMemcpyDeviceToHost);
//Free device matrices
cudaFree(gpu_A);
cudaFree(gpu_B);
cudaFree(gpu_C);
}
}
A função vecmul() é o kernel e a função maxmul() é a helper. Sua função é alocar memória na GPU, copiar os parâmetros, invocar o kernel, e copiar o resultado. Os valores são passados por referência.
O programa maxmul.go invoca a função helper e exibe o resultado:
package main
/*
void maxmul(float *A, float* B, float *C, int size);
#cgo LDFLAGS: -L. -L./ -lmaxmul
*/
import "C"
import "fmt"
func Maxmul(a []C.float, b []C.float, c []C.float, size int) {
C.maxmul(&a[0], &b[0], &c[0], C.int(size))
}
func main() {
//in := []C.float{1.23, 4.56}
//C.test(&in[0]) // C 1.230000 4.560000
a := []C.float{-1,2,4,0,5,3,6,2,1}
b := []C.float{3,0,2,3,4,5,4,7,2}
var c []C.float = make([]C.float, 9)
Maxmul(a,b,c,3)
fmt.Println(c)
}
Antes da importação do pacote C, que permite invocar funções externas em código C puro (extern C), eu passo a configuração do cgo, indicando o protótipo da função C, o caminho para a lib e o seu nome.
Tive que criar uma wrapper function no código Go para invocar a função externa, de modo a facilitar as coisas. Ela simplesmente passa a referência para os arrays (o endereço da primeira posição) e o tamanho da matriz (neste caso 3x3 = 9). Em CUDA trabalhamos com matrizes achatadas.
Usei o tipo C.float para criar slices contendo minhas matrizes (transformadas em vetores). Depois, invoquei a função. Note que passei o tamanho de cada linha (ou coluna).
Para compilar o código C use o comando:
nvcc --ptxas-options=-v --compiler-options '-fPIC' -o libmaxmul.so --shared maxmul.cu
Você precisa ter o CUDA e o driver Nvidia instalados!
Depois, é só rodar o código Go com o comando:
LD_LIBRARY_PATH=${PWD}${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}} go run maxmul.go
...
[19 36 16 27 41 31 28 15 24]
Você deve adicionar LD_LIBRARY_PATH = $ {PWD} $ {LD_LIBRARY_PATH: +: $ {LD_LIBRARY_PATH}} ao colocar a biblioteca em um local diferente do já definido em LD_LIBRARY_PATH. Você pode encontrar mais informações sobre esta variável de ambiente [aqui] (https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html#post-installation-actions)
E este é o resultado da multiplicação das matrizes!