Programação para GPU

Aleardo Manacero Jr.

Introdução

O uso de GPUs em computação de alto desempenho começou com o desenvolvimento de APIs que permitissem a programação das placas gráficas em alto nível

Isso significa abandonar bibliotecas gráficas, como DirectX e usar primitivas apropriadas à paralelização de aplicações

Introdução

Três bibliotecas para programação de GPU:

CUDA

OpenCL

DirectCompute

Esse processo, infelizmente, depende da arquitetura de cada máquina

Arquitetura de GPUs

GPUs são dispositivos originalmente criados para acelerar o processamento gráfico

Sua arquitetura é massivamente paralela

Os elementos de processamento são direcionados para operações aritméticas

O alto grau de paralelismo dá seu desempenho

Arquitetura de GPUs

ATI Radeon 5870

Arquitetura de GPUs

ATI Radeon 5870

Arquitetura de GPUs

Nvidia Kepler

Arquitetura de GPUs

Nvidia Kepler

Programação

As diferenças de arquitetura tornam a programação sensível ao contexto

Um programa otimizado para uma GPU pode não funcionar bem em outra

O que fazer então?

Programação

Usar técnicas de particionamento do problema adequadas para a GPU a ser usada

Usar a biblioteca para paralelização adequada para a GPU a ser usada

openCL

Programação

OpenCL (Open Computing Language) é uma API para programação de GPUs (ou outros dispositivos) e sua interação com a CPU

Framework de três partes (linguagem, plataforma e execução)

Programação

Linguagem baseada em C99, com extensões e restrições adicionais

API de Plataforma com rotinas para tratar o sistema e seus recursos

API de execução gerencia objetos criados na execução de processos openCL

Kernel

Programas são divididos em dois componentes:

Host, que executa na CPU e envia kernels para as GPUs

Kernel, que executa nas GPUs, de modo paralelo (dados ou tarefas)

Kernel

Para a paralelização de kernel é preciso definir o número de dimensões do objeto a ser tratado

Cada elemento no objeto é chamado de work-item

O número de kernels a ser criado depende da dimensão global do objeto

Kernel

Pode haver aglomeração de work-itens em um work-group

Kernel

Pode haver aglomeração de work-itens em um work-group

Exemplo (paralelismo de dados)

SEQUENCIAL

void square(int n, const float *a,

float *result)

{ int i;

for (i=0; i<n; i++)

result[i] = a[i] * a[i];

}

PARALELO

kernel dp_square (const float *a,

float *result)

{ int id = get_global_id(0);

result[id] = a[id] * a[id];

}

// dp_square executes over “n” work-items

Host

Gerencia os recursos openCL, ou seja, gerencia:

Dispositivos

Rotinas (código) de kernel

Kernel

Memória

Memória

É gerenciada explicitamente pelo programa

Cuida das movimentações entre host e kernel

Se divide em quatro níveis dentro da GPU

Memória

Na GPU temos:

Global – R/W para todos itens e groups

Constante – R/W para host e R para itens

Local – R/W para itens no mesmo group

Privada – R/W para um item apenas

Memória

Execução de um programa

Envolve várias etapas:Buscar dispositivos no sistema

Criar contexto para associar dispositivos

Criar programas para execução nos dispositivos

Selecionar kernels para execução

Criar objetos de memória (host e dispositivo)

Transferir dados para dispositivos

Submeter kernels para fila de comandos

Copiar resultados para o host

Execução de um programa

Identificação das plataformas disponíveis

cl_platform_id platforms;cl_uint num_platforms;

// query for 1 available platformcl_int err = clGetPlatformIDs (

1, // the number of entries that can be added to platforms&platforms, // list of OpenCL found&num_platforms); // the number of OpenCL platforms found

// platforms are NVIDIA, INTEL, AMD, ...

Execução de um programa

Identificação dos dispositivos disponíveis

cl_device_id device_id;cl_uint num_of_devices;cl_int err;err = clGetDeviceIDs (platform_id, // platform_id retrieved from clGetPlatformIDsCL_DEVICE_TYPE_GPU, // device type to search for1, // number of id add to device_id list&device_id, // list of device ids&num_of_devices // number of compute devices found);

TYPE pode ser CPU, GPU, ACCELERATOR, DEFAULT, ALL

Execução de um programa

Criação de contexto

cl_context context; // context properties list - must be ended with 0properties[0]= CL_CONTEXT_PLATFORM; // tells platform to useproperties[1]= (cl_context_properties) platform_id;properties[2]= 0;context = clCreateContext (

properties, // list of context properties1, // num of devices in the device_id list&device_id, // the device id listNULL, // pointer to the error callback function (if required)NULL, // the argument data to pass to the callback function&err // the return code

);

Execução de um programa

Criação da fila de comandos

cl_command_queue command_queue;command_queue = clCreateCommandQueue (

context, // a valid contextdevice_id, // a valid device associated with the context0, // properties for the queue - not used here&err // the return code

);

Campo properties (0 no exemplo) pode ser:CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE CL_QUEUE_PROFILING_ENABLE

Execução de um programa

Criação de um programa

const char *ProgramSource ="__kernel void hello (__global float *input, __global float *output)\n"\"{\n"\" size_t id = get_global_id (0);\n"\" output[id] = input[id] * input[id];\n"\"}\n";

Neste exemplo o programa aparece na forma de string a ser lida pelo código do host

Execução de um programa

Criação de um programa

cl_program program;program = clCreateProgramWithSource (

Context, // a valid context1, // the number strings in the next parameter(const char **) &ProgramSource, // the array of stringsNULL, // the length of each string or can be NULL terminated&err // the error return code

);

Se o programa já for um executável trocar por clCreateProgramWithBinary

Execução de um programa

Criação do executável

err = clBuildProgram (program, // a valid program object0, // number of devices in the device listNULL, // device list – NULL means for all devicesNULL, // a string of build optionsNULL, // callback function when executable has been builtNULL // data arguments for the callback function

);

Execução de um programa

Criação do executável

err = clBuildProgram (program, // a valid program object0, // number of devices in the device listNULL, // device list – NULL means for all devicesNULL, // a string of build optionsNULL, // callback function when executable has been builtNULL // data arguments for the callback function

);

clGetBuildProgramInfo() permite recuperar o log da função acima, para eventuais erros de compilação

Execução de um programa

Criação de objetos kernel

cl_kernel kernel;

kernel = clCreateKernel(program, // a valid program object successfully built"hello", // the name of the kernel declared with __kernel&err // error return code

);

clCreateKernel retorna um valor diferente de zero se for bem sucedida ou zero se falhar (tipo de erro é retornado em “err”)

Execução de um programa

Erros podem ser:

CL_INVALID_PROGRAM — O programa não é um objeto válido

CL_INVALID_PROGRAM_EXECUTABLE — O programa não contém um executável construído com sucesso

CL_INVALID_KERNEL_NAME — Nome do kernel não foi localizado no programa objeto

CL_INVALID_VALUE — Nome do kernel é NULL.

CL_OUT_OF_HOST_MEMORY — Host não foi capaz de alocar recursos OpenCL

Execução de um programa

Definição de parâmetros do kernel

err = clSetKernelArg (kernel, // valid kernel object0, // the specific argument index of a kernelsizeof(cl_mem), // the size of the argument data&input_data // a pointer of data used as the argument

);

clGetKernelInfo() retorna informações sobre os parâmetros de um kernel

Execução de um programa

Colocando um kernel para execução

err = clEnqueueNDRangeKernel (command_queue, // valid command queuekernel, // valid kernel object1, // the work problem dimensionsNULL, // reserved for future revision - must be NULL&global, // work-items for each dimensionNULL, // work-group size for each dimension0, // number of event in the event listNULL, // list of events that needs to complete before this executesNULL // event object to return on completion

);

Execução de um programa

Para o comando anterior é preciso definir alguns parâmetros:

O número de dimensões pode ser 1, 2 ou 3

work-itens por dimensão precisa ser definida, como

size_t global[2]={512,512};

work-group por dimensão deve ser definido, como

size_t local[2]={8,8};

Execução de um programa

Terminando a execução (liberando recursos)

clReleaseMemObject(input);

clReleaseMemObject(output);

clReleaseProgram(program);

clReleaseKernel(kernel);

clReleaseCommandQueue(command_queue);

clReleaseContext(context);

Vetores em openCL

Vetores são definidos de forma explícita, como por exemplo int8 alfa;

Vetores em openCL

O número de elementos pode ser 2, 4, 8 ou 16

Os elementos podem ser indexados por valores hexadecimais ou letras (até 4 elem.)

float16 x;x.sa = 1.0f;

float4 c;c.x = 1.0f; c.y = 1.0f; c.z = 1.0f; c.w = 1.0f;

Vetores em openCL

O acesso ainda pode ser feito por grupos, como:

float4 f = (float4) (1.0f, 2.0f, 3.0f, 4.0f);float2 low, high;float2 o, e;

low = f.lo; // returns f.xy (1.0f, 2.0f)high = f.hi; // returns f.zw (3.0f, 4.0f)

o = f.odd; // returns f.yw (2.0f, 4.0f)e = f.even; // returns f.xz (1.0f, 3.0f)

Vetores em openCL

Casting é feito por funções de conversão, apesar de tipos escalares também aceitarem casting clássico

int4 i;float4 f = convert_float4(i); // int4 vector to float4

float f;int i = convert_int(f); // float scaler to an integer scaler

Exemplo em openCL

Multiplicação convencional de matrizes

Alguns trechos foram suprimidos (a geração dos conteúdos das matrizes, p.ex.)

O código do kernel está em arquivo lido pelo programa do host

As matrizes são tratadas como vetores

Host code openCL

#include <stdio.h> + <stdlib.h> + <string.h> + <math.h>#include <fcntl.h>#include <unistd.h>#include <sys/types.h>#include <sys/stat.h>#include <stdbool.h>#include <CL/cl.h> // openCL headers

#define WA 1024#define HA 1024#define WB 1024

#define HB WA#define WC WB#define HC HA

Host code openCL

long LoadOpenCLKernel(char const* path, char **buf){ FILE *fp; size_t fsz; long off_end; int rc;

/* Open the file */ fp = fopen(path, "r"); if ( NULL == fp ) { return -1L; }

/* Seek to the end of the file */ rc = fseek(fp, 0L, SEEK_END); if ( 0 != rc ) { return -1L; }

Host code openCL

/* Byte offset to the end of the file (size) */ if ( 0 > (off_end = ftell(fp)) ) { return -1L; } fsz = (size_t)off_end;

/* Allocate a buffer to hold the whole file */ *buf = (char *) malloc( fsz+1); if ( NULL == *buf ) { return -1L; }

/* Rewind file pointer to start of file */ rewind(fp);

Host code openCL

/* Slurp file into buffer */ if( fsz != fread(*buf, 1, fsz, fp) ) { free(*buf); return -1L; }

/* Close the file */ if( EOF == fclose(fp) ) { free(*buf); return -1L; }

/* Make sure the buffer is NUL-terminated, just in case */ (*buf)[fsz] = '\0';

/* Return the file size */ return (long)fsz;}

Host code openCL

int main(int argc, char** argv){int err; // error code returned from api calls cl_device_id device_id; // compute device id cl_context context; // compute context cl_command_queue commands; // compute command queue cl_program program; // compute program cl_kernel kernel; // compute kernel

// OpenCL device memory for matrices cl_mem d_A; cl_mem d_B; cl_mem d_C;

Host code openCL

//Allocate host memory for matrices A and B unsigned int size_A = WA * HA; unsigned int mem_size_A = sizeof(float) * size_A; float* h_A = (float*) malloc(mem_size_A); unsigned int size_B = WB * HB; unsigned int mem_size_B = sizeof(float) * size_B; float* h_B = (float*) malloc(mem_size_B);

//Initialize host memory randomMemInit(h_A, size_A); randomMemInit(h_B, size_B); //Allocate host memory for the result C unsigned int size_C = WC * HC; unsigned int mem_size_C = sizeof(float) * size_C; float* h_C = (float*) malloc(mem_size_C);

Host code openCL

printf("Initializing OpenCL device...\n");

cl_uint dev_cnt = 0; clGetPlatformIDs(0, 0, &dev_cnt);

cl_platform_id platform_ids[100]; clGetPlatformIDs(dev_cnt, platform_ids, NULL);

// Connect to a compute device int gpu = 1; err = clGetDeviceIDs(platform_ids[0], gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL); if (err != CL_SUCCESS) { printf("Error: Failed to create a device group!\n"); return EXIT_FAILURE; }

Host code openCL

// Create a compute context context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); if (!context) { printf("Error: Failed to create a compute context!\n"); return EXIT_FAILURE; }

// Create a command commands commands = clCreateCommandQueue(context, device_id, 0, &err); if (!commands) { printf("Error: Failed to create a command commands!\n"); return EXIT_FAILURE; }

Host code openCL

// Create the compute program from the source file char *KernelSource; long lFileSize;

lFileSize = LoadOpenCLKernel("matrixmul_kernel.cl", &KernelSource); if( lFileSize < 0L ) { perror("File read failed"); return 1; }

program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err); if (!program) { printf("Error: Failed to create compute program!\n"); return EXIT_FAILURE; }

Host code openCL

// Build the program executable

err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); if (err != CL_SUCCESS) { size_t len; char buffer[2048]; printf("Error: Failed to build program executable!\n"); clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len); printf("%s\n", buffer); exit(1); }

// Create the compute kernel in the program we wish to run

kernel = clCreateKernel(program, "matrixMul", &err); if (!kernel || err != CL_SUCCESS) { printf("Error: Failed to create compute kernel!\n"); exit(1); }

Host code openCL

// Create the input and output arrays in device memory for our calculation d_C = clCreateBuffer(context, CL_MEM_READ_WRITE, mem_size_A, NULL, &err); d_A = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, mem_size_A, h_A, &err); d_B = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, mem_size_B, h_B, &err);

if (!d_A || !d_B || !d_C) { printf("Error: Failed to allocate device memory!\n"); exit(1); }

Host code openCL

printf("Running matrix multiplication for matrices A (%dx%d) and B (%dx%d) ...\n", WA,HA,WB,HB);

//Launch OpenCL kernel size_t localWorkSize[2], globalWorkSize[2];

int wA = WA; int wC = WC; err = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_C); err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_A); err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&d_B); err |= clSetKernelArg(kernel, 3, sizeof(int), (void *)&wA); err |= clSetKernelArg(kernel, 4, sizeof(int), (void *)&wC);

if (err != CL_SUCCESS) { printf("Error: Failed to set kernel arguments! %d\n", err); exit(1); }

Host code openCL

localWorkSize[0] = 16; localWorkSize[1] = 16; globalWorkSize[0] = 1024; globalWorkSize[1] = 1024; err = clEnqueueNDRangeKernel(commands, kernel, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);

if (err != CL_SUCCESS) { printf("Error: Failed to execute kernel! %d\n", err); exit(1); } //Retrieve result from device err = clEnqueueReadBuffer(commands, d_C, CL_TRUE, 0, mem_size_C, h_C, 0, NULL, NULL);

Host code openCL

if (err != CL_SUCCESS) { printf("Error: Failed to read output array! %d\n", err); exit(1); }

printf("Matrix multiplication completed...\n");

//Shutdown and cleanup free(h_A); free(h_B); free(h_C);

clReleaseMemObject(d_A); clReleaseMemObject(d_C); clReleaseMemObject(d_B);

clReleaseProgram(program); clReleaseKernel(kernel); clReleaseCommandQueue(commands); clReleaseContext(context);

return 0;}

Kernel em openCL

// Arquivo matrixmul_kernel.cl

// OpenCL Kernel__kernel void matrixMul(__global float* C, __global float* A, __global float* B, int wA, int wB){ int tx = get_global_id(0); int ty = get_global_id(1); // value stores the element that is computed by the thread float value = 0;

Kernel em openCL

for (int k = 0; k < wA; ++k) { float elementA = A[ty * wA + k]; float elementB = B[k * wB + tx]; value += elementA * elementB; } // Write the matrix to device memory each // thread writes one element C[ty * wA + tx] = value;}

CUDA

Introdução

Criada para programação de GPUs da Nvidia

Programação mais simples do que openCL

Problema é ser restrita ao fabricante

Usaremos material da própria Nvidia