Programação CUDA em C#

A tecnologia CUDA, acrónimo para Compute Unified Device Architecture, foi desenvolvida pelo fabricante de placas gráficas NVIDIA e consiste numa plataforma e modelo de programação em C/C++ vocacionada para o aumento do desempenho das aplicações com o auxílio dos GPU (Graphics Processing Unit). Dada a especificidade dos dispositivos gráficos, a sua arquitectura viabiliza a realização paralela de operações matemáticas com um desempenho superior aos CPUs (Central Processing Unit) com custos reduzidos. Esta característica tem disseminado o recurso a este tipo de dispositivos em HPC (High Performance Computing) o que, por si só, é suficiente para considerar o domínio da tecnologia como uma mais-valia. A alternativa ao CUDA, definida pelo consórcio Khronos Group, é proporcionada pelo padrão livre de computação paralela em sistemas heterogéneos OpenCL (Open Computing Language). Este padrão é transversal a um número significativo e crescente de fabricantes de dispositivos tais como a Intel, IBM, AMD ou ARM, incluindo-se mesmo nesta lista a NVIDIA.

Não pretendo aqui justificar a utilização de uma tecnologia em detrimento de outra. Pretendo, por outro lado, descrever como é possível utilizar a tecnologia CUDA com linguagens alternativas como é o caso do C# ou do JAVA. Escolhi descrever CUDA devido ao facto de já ter efectuado algumas digressões neste sentido.

Princípios básicos

É possível encontrar no Cuda Toolkit Documentation toda a informação necessária para desenvolver aplicações CUDA nas linguagens C e C++. Aí encontram-se também definidos os passos para a instalação e configuração. Será suposto, nos exemplos que se seguem, que tantos os drivers da NVIDIA como o respectivo CUDA Toolkit se encontram propriamente instalados e configurados. Supôr-se-á ainda que o sistem onde os programas são executados contêm dispositivos com suporte CUDA.

Começamos por considerar o código de exemplo

// Device code
extern "C" __global__ void VecAdd(float* A, float* B, float* C, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// Host code
int main()
{
int i;
int N = 3;
size_t size = N * sizeof(float);

// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
float* h_C = (float*)malloc(size);

// Initialize input vectors
for(i = 0; i &amp;amp;lt; N; ++i)
{
h_A[i] = (float)i;
h_B[i] = (float)(N + i);
}

// Allocate vectors in device memory
float* d_A;
cudaMalloc(d_A, size);
float* d_B;
cudaMalloc(d_B, size);
float* d_C;
cudaMalloc(d_C, size);

// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

// Invoke kernel
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<<blocksPerGrid, threadsPerBloc>>>(d_A, d_B, d_C, N);

// Copy result from device memory to host memory
// h_C contains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);

// Print result
for(i = 0; i < N; ++i)
{
printf("%f", h_A[i]);
}

printf("\n");

for(i = 0; i < N; ++i)
{
printf("%f", h_B[i]);
}

printf("\n");
for(i = 0; i < N; ++i)
{
printf("%f", h_C[i]);
}

// Free host memory
if(h_A != NULL) free(h_A);
if(h_B != NULL) free(h_B);
if(h_C != NULL) free(h_C);
}

Tata-se da codificação de um programa que recorre a uma placa gráfica para efectuar a soma de dois vectores. Se supusermos que este código se encontra definido no ficheiro AddVector.cu (com extensão CUDA), a sua compilação num programa executável pode ser conseguida com o auxílio do comando:

nvcc -arch=sm_20 -rdc=true AddVector.cu -o AddVector -lcudadevrt -ccbin “C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin\x86_amd64”

Assumimos aqui que o Visual Studio 2012 se encontra instalado no computador, uma vez que versões do compilador não são suportadas pelo programa nvcc. De facto, a directoria especificada na opção -ccbin deverá conter qualquer um dos compiladores suportados. A opção -arch permite indiciar ao compilador que o código a ser gerado deverá ser executado por dispositivos que suportem a arquitectura 2.0.

A execução do comando anterior permite-nos criar o programa AddVector.exe que imprime o seguinte resultado:

0.000000 1.000000 2.000000
3.000000 4.000000 5.000000
3.000000 5.000000 7.000000

De um modo resumido, as funções marcadas com o qualificador __global__ são executadas na gráfica mas chamadas no anfitrião. As funções marcadas como __device__ são chamadas e executadas na gráfica. As restantes funções são de anfitrião, cujo qualificador é __host__. No código acima, percebemos a definição da função

VecAdd(float*, float*, float*,int)

que é chamada no código de anfitrião na função main,

VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C)

As funções cudaMalloccudaMemcpycudaFree são definidas no CUDA Runtime API que constitui uma interface CUDA para C/C++.

O compilador nvcc

A abordagem anterior não se presta ao desenvolvimento de aplicações CUDA baseadas em outras linguagens de programação por razões óbvias, entre as quais se destaca a utilização da Runtime API. Assim, torna-se necessário perceber como o compilador gera o executável final.

Na secção The CUDA Compilation Trajectory do manual são descritas as fases da compilação de uma aplicação CUDA. O programa de entrada é pré-processado para compilação em binários CUDA (cubin) ou código intermédio (PTX) de dispositivo, os quais são colocados num ficheiro binário fatbinary. O programa de entrada é processado para compilação de anfitrião com o auxílio do compilador C/C++ especificado na linha de comandos. Neste processo, o binário fatbinay é imbuído no código de anfitrião onde as extensões CUDA são tranformadas em construções C++. O resultado final é compilado num objecto de anfitrião. Quando o código de dispositivo é lançado, o fatbinary é inpeccionado de modo a que o programa de anfitrião obtenha a imagem correcta para o GPU actual. O utilitário de compilação CUDA, como descrito na secção Using Separate Compilation in CUDA, pode ser configurado de modo a separar todo o processo de compilação de vários ficheiros CUDA. Esta funcionalidade encontra-se disponível desde a versão 5.0.

Cada etapa da compilação CUDA gera, numa directoria temporária, cada um dos ficheiros resultantes, os quais são eliminados imediatamente antes da sua conclusão. No entanto, como descrito na secção Keeping intermediate phase files o nvcc pode ser configurado por intermédio de uma opção de modo a que todos os ficheiros gerados sejam mantidos na directoria de compilação. Este tipo de compilação pode ser conseguido com o auxílio da execução da seguinte linha de comandos:

nvcc -arch=sm_20 -rdc=true AddVector.cu -o AddVector -lcudadevrt -ccbin “C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin\x86_amd64” -keep

É possível identificar três ficheiros, nomeadamente AddVector.ptxAddVector.sm_20.cubinAddVector.fatbin. O primeiro ficheiro, do tipo ptx consiste num conjunto textual de instruções muito semelhante ao assembly para processadores de CPU. O código contido neste ficheiro é orientado para um determinado conjunto de arquitecturas virtuais. Numa outra fase da compilação do nvcc, são gerados os binários cubin que são suportados por uma arquitectura real (e que poderão não funcionar com outras arquitecturas). No entanto, a arquitectura real deverá constituir uma implementação da arquitectura virtual de modo a que a compilação seja bem sucedida. Este processo permite mitigar o efeito da compatibilidade das aplicações face à evolução das GPU. De facto, o driver CUDA permite a compilação de ficheiros ptx em tempo de execução durante o qual já é conhecida a arquitectura real da GPU. Os binários fatbinary contêm os códigos de várias arquitecturas reais que implementam uma determinada arquitectura virtual.

O CUDA driver API

É natural assumirmos que o cubin imbuído no distribuível é carregado no GPU aquando das chamadas do kernel. Tal carregamento é conseguido com o auxílio de funções de mais baixo nível descritas no CUDA driver API. No sistema operativo Windows, essas funções são disponibilizadas pela dll nvcuda.dll instalada na pasta System32.

É possível mapear a interface disponibilizada pela driver API em funções C# com o auxílio do atributo DllImport. De facto, tal mapeamento pode ser encontrado no projecto matutils. A título de exemplo, seja considerado seguinte código:


static void Main(string[] args)
 {
 var cudaResult = CudaApi.CudaInit(0);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

// Obtém o número de dispositivos disponíveis
 var deviceCount = default(int);
 cudaResult = CudaApi.CudaDeviceGetCount(ref deviceCount);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

 for (int i = 0; < deviceCount; ++i)
 {
 var currentDevice = default(int);
 cudaResult = CudaApi.CudaDeviceGet(ref currentDevice, i);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

 var deviceName = new StringBuilder();
 cudaResult = CudaApi.CudaDeviceGetName(deviceName, 64, i);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

 var major = default(int);
 var minor = default(int);
 cudaResult = CudaApi.CudaDeviceGetAttribute(
 ref major,
 ECudaDeviceAttr.ComputeCapabilityMajor,
 i);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

cudaResult = CudaApi.CudaDeviceGetAttribute(
 ref minor,
 ECudaDeviceAttr.ComputeCapabilityMinor,
 i);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("CUDA function failed.");
 }

Console.WriteLine(
 "GPU: {0}; Compute Capability: {1}.{2}",
 deviceName,
 major,
 minor);
 }

}

Constata-se facilmente que a primeira etapa consiste na inicialização do driver. A execução da função CuInit, mapeada no matutils como CudaInit, é obrigatória sempre que se pretenda recorrer às GPU em CUDA. O seguimento do código permite obter o número de dispositivos CUDA conectados ao anfitrião e obter, relativamente a cada um deles, o respectivo nome e poder de computação. A compilação do AddVector.cu com a linha de comandos

nvcc -arch=compute_30 -code=sm_30 -rdc=true AddVector.cu -lcudadevrt -ccbin “C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin\x86_amd64” -cubin -o AddVector.sm_30.cubin

permite gerar um ficheiro cubin que pode ser carregado e executado numa GPU com poder computacional 3.0. O código que se segue permite carregar o módulo que contém a função VecAdd compilado anteriormente e executá-lo com o auxílio do lançamento de um kernel.


static void Main(string[] args)
 {
 // Inicializa CUDA e avalia os dispositivos existentes
 var cudaResult = CudaApi.CudaInit(0);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Obtém o primeiro dispositivo
 var device = default(int);
 cudaResult = CudaApi.CudaDeviceGet(ref device, 0);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// O contexto é automaticamente colocado
// como corrente para a linha de fluxo actual
 var context = default(SCudaContext);
 cudaResult = CudaApi.CudaCtxCreate(
ref context,
ECudaContextFlags.SchedAuto, device);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Carrega o módulo no contexto actual
 var module = default(SCudaModule);
 cudaResult = CudaApi.CudaModuleLoad(
ref module,
"Data\\AddVector.sm_30.cubin");
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Obtém a função a ser chamada
 var cudaFunc = default(SCudaFunction);
 cudaResult = CudaApi.CudaModuleGetFunction(
 ref cudaFunc,
 module,
 "VecAdd");
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

var elemensNum = 10;

//var start = 0;
 var firstVector = new int[elemensNum];
 var secondVector = new int[elemensNum];
 var result = new int[elemensNum];
 for (int i = 0; i < elemensNum; ++i)
 {
 firstVector[i] = i + 1;
 secondVector[i] = elemensNum - i;
 }

// Reserva o primeiro vector
 var firstCudaVector = default(SCudaDevicePtr);
 cudaResult = CudaApi.CudaMemAlloc(
 ref firstCudaVector,
 Marshal.SizeOf(typeof(int)) * elemensNum);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Reserva o segundo vector
 var secondCudaVector = default(SCudaDevicePtr);
 cudaResult = CudaApi.CudaMemAlloc(
 ref secondCudaVector,
 Marshal.SizeOf(typeof(int)) * elemensNum);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Reserva o terceiro vector
 var resultCudaVector = default(SCudaDevicePtr);
 cudaResult = CudaApi.CudaMemAlloc(
 ref resultCudaVector,
 Marshal.SizeOf(typeof(int)) * elemensNum);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

var cudaSize = default(SCudaDevicePtr);
 cudaResult = CudaApi.CudaMemAlloc(
 ref cudaSize,
 Marshal.SizeOf(typeof(int)));
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Efectua a cópia do primeiro vector para o dispositivo
 var handle = GCHandle.Alloc(
firstVector,
GCHandleType.Pinned);
 var size = Marshal.SizeOf(typeof(int));
 var hostPtr = handle.AddrOfPinnedObject();

cudaResult = CudaApi.CudaMemcpyHtoD(
 firstCudaVector,
 hostPtr,
 elemensNum * size);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

handle.Free();

// Efectua a cópia do segundo vector para o dispositivo
 handle = GCHandle.Alloc(
secondVector,
GCHandleType.Pinned);
 hostPtr = handle.AddrOfPinnedObject();

cudaResult = CudaApi.CudaMemcpyHtoD(
 secondCudaVector,
 hostPtr,
 elemensNum * size);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

handle.Free();

var vectorSizePtr = Marshal.AllocHGlobal(
Marshal.SizeOf(typeof(int)));
 cudaResult = CudaApi.CudaMemcpyHtoD(
 cudaSize,
 hostPtr,
 size);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

Marshal.FreeHGlobal(vectorSizePtr);

// Reserva espaço para o vector de argumentos do kernel
 var managedPtrArray = new IntPtr[4];
 var ptrSize = Marshal.SizeOf(typeof(IntPtr));
 var unmanagedArrayPtr = Marshal.AllocHGlobal(ptrSize * 3);

// Procede à criação dos objectos em código não gerido
 var managedElementPtr = Marshal.AllocHGlobal(Marshal.SizeOf(
typeof(SCudaDevicePtr)));
 managedPtrArray[0] = managedElementPtr;
 Marshal.StructureToPtr(
firstCudaVector,
 managedElementPtr,
false);
 Marshal.WriteIntPtr(unmanagedArrayPtr, 0, managedElementPtr);

managedElementPtr = Marshal.AllocHGlobal(
Marshal.SizeOf(typeof(SCudaDevicePtr)));
 managedPtrArray[1] = managedElementPtr;
 Marshal.StructureToPtr(
secondCudaVector,
managedElementPtr,
false);
 Marshal.WriteIntPtr(
unmanagedArrayPtr,
ptrSize,
managedElementPtr);

managedElementPtr = Marshal.AllocHGlobal(
Marshal.SizeOf(typeof(SCudaDevicePtr)));
 managedPtrArray[2] = managedElementPtr;
 Marshal.StructureToPtr(
resultCudaVector,
managedElementPtr,
false);
 Marshal.WriteIntPtr(
unmanagedArrayPtr,
2 * ptrSize,
managedElementPtr);

managedElementPtr = Marshal.AllocHGlobal(
Marshal.SizeOf(typeof(SCudaDevicePtr)));
 managedPtrArray[3] = managedElementPtr;
 Marshal.StructureToPtr(
cudaSize,
 managedElementPtr,
false);
 Marshal.WriteIntPtr(
unmanagedArrayPtr,
3 * ptrSize,
managedElementPtr);

// Realiza a chamada
 cudaResult = CudaApi.CudaLaunchKernel(
 cudaFunc,
 (uint)elemensNum,
 1,
 1,
 1,
 1,
 1,
 0,
 new SCudaStream(),
 unmanagedArrayPtr,
 IntPtr.Zero);

cudaResult = CudaApi.CudaCtxSynchronize();

// Liberta o conjunto de argumentos reservado
 Marshal.FreeHGlobal(unmanagedArrayPtr);
 for (int i = 0; i &amp;amp;lt; 4; ++i)
 {
 var current = managedPtrArray[i];
 Marshal.FreeHGlobal(current);
 }

// Copia de volta o terceiro vector para o anfitrião
 handle = GCHandle.Alloc(result, GCHandleType.Pinned);
 hostPtr = handle.AddrOfPinnedObject();
 cudaResult = CudaApi.CudaMemcpyDtoH(
 hostPtr,
 resultCudaVector,
 size * elemensNum);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

handle.Free();

// Imprime o conteúdo dos vectores
 for (int i = 0; i < elemensNum; ++i)
 {
 Console.Write("{0} ", firstVector[i]);
 }

Console.WriteLine();

for (int i = 0; i < elemensNum; ++i)
 {
 Console.Write("{0} ", secondVector[i]);
 }

Console.WriteLine();

for (int i = 0; i < elemensNum; ++i)
 {
 Console.Write("{0} ", result[i]);
 }

// Liberta o primeiro vector do GPU
 cudaResult = CudaApi.CudaMemFree(firstCudaVector);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Liberta o segundo vector do GPU
 cudaResult = CudaApi.CudaMemFree(secondCudaVector);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Liberta o vector do resultado do GPU
 cudaResult = CudaApi.CudaMemFree(resultCudaVector);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Liberta o espaço reservado para conter o número
// de elementos de cada vector
 cudaResult = CudaApi.CudaMemFree(cudaSize);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Remove o módulo do contexto actual
 cudaResult = CudaApi.CudaModuleUnload(module);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }

// Descarta o contexto
 cudaResult = CudaApi.CudaCtxDestroy(context);
 if (cudaResult != ECudaResult.CudaSuccess)
 {
 throw new Exception("A CUDA error has occurred.");
 }
 }

Dada a pretensão de utilizar CUDA, é necessário, em primeiro lugar, inicializar o driver CUDA. De seguida, é obtido o primeiro dispositivo registado cujo valor ordinal é 0. Um contexto é muito semelhante a um processo de anfitrião onde, por exemplo, cada conjunto de variáveis tem um espaço de endereçamento próprio. As funções são, portanto, carregadas e executadas no âmbito de um processo. Todos os recursos reservados num determinado contexto são libertados aquando da sua destruição. Os passos que se seguem à criação do contexto são o carregamento do módulo que contém a função que pretendemos executar que, neste caso, é a VecAdd, a obtenção do seu apontador para chamada de execução a partir do anfitrião e o lançamento do kernel respectivo. Como as funções do driver executam no exterior da máquina virtual do .NET, é necessária uma interligação entre processos .NET e processos de código não gerido. Não serão aqui discutidas as  formas de transferir dados entre código gerido e não gerido.

Após a execução do kernel é sincronizado o contexto do dispositivo com o processo de anfitrião. Apesar da destruição do contexto libertar os recursos reservados, talvez constitua uma boa prática implementar a libertação dos recursos do contexto de modo a ser possível reutilizá-lo em computações futuras.

 

Sobre Sérgio O. Marques

Licenciado em Física/Matemática Aplicada (Astronomia) pela Faculdade de Ciências da Universidade do Porto e Mestre em Matemática Aplicada pela mesma instituição, desenvolvo trabalho no PTC (Porto Technical Centre) - Yazaki como Administrador de bases-de-dados. Dentro o meu leque de interesses encontram-se todos os temas afins às disciplinas de Matemática, Física e Astronomia. Porém, como entusiasta, interesso-me por temas relacionados com electrónica, poesia, música e fotografia.
Esta entrada foi publicada em Computadores e Internet com as etiquetas , , , . ligação permanente.

Deixe uma Resposta

Preencha os seus detalhes abaixo ou clique num ícone para iniciar sessão:

Logótipo da WordPress.com

Está a comentar usando a sua conta WordPress.com Terminar Sessão / Alterar )

Imagem do Twitter

Está a comentar usando a sua conta Twitter Terminar Sessão / Alterar )

Facebook photo

Está a comentar usando a sua conta Facebook Terminar Sessão / Alterar )

Google+ photo

Está a comentar usando a sua conta Google+ Terminar Sessão / Alterar )

Connecting to %s