W kolejnej części wprowadzenia do OpenCL wykonamy i uruchomimy pierwszy program na GPU. Zadaniem programu będzie sumowanie dwóch wektorów.
Przykładowy kod w C++ dla takiej operacji może wyglądać tak:
const int vectorSize = 10;
int a[vectorSize];
int b[vectorSize];
int c[vectorSize];
for (int i = 0; i < vectorSize; ++i)
{
c[i] = a[i] + b[i];
}
Pobieramy kod z poprzedniej części kursu.
Na początku utworzymy tablice przechowujące wektory:
size_t vectorNumber = 32;
cl_int* a = new cl_int[vectorSize];
cl_int* b = new cl_int[vectorSize];
cl_int* c = new cl_int[vectorSize];
randomizeArray(a, vectorSize);
randomizeArray(b, vectorSize);
Zmienna „vectorNumber” przechowuje rozmiar wektora. Funkcja „randomizeArray” wypełnia tablice losowymi wartościami. Jej kod zostanie przedstawiony później.
Pierwszą czynnością, którą musimy wykonać jest utworzenie kontekstu, który umożliwi nam wykonywanie poleceń na danym urządzeniu.
cl_context context = clCreateContext(0, deviceNumber, deviceIds, NULL, NULL, NULL);
if (NULL == context)
{
std::cout << "Failed to create OpenCL context." << std::endl;
}
Następnie tworzymy kolejkę poleceń na wybranym przez nas urządzeniu:
cl_command_queue commandQueue = clCreateCommandQueue(context, deviceIds[0], 0, &error);
Alokujemy pamięć na karcie graficznej dla naszych wektorów. Tworzymy dwa bufory tylko do odczytu, które będą przechowywać wektory A i B oraz jeden tylko do zapisu, do którego zostanie zapisany wynik sumowania:
cl_mem bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int) * vectorSize, NULL, &error);
cl_mem bufferB = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int) * vectorSize, NULL, &error);
cl_mem bufferC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_int) * vectorSize, NULL, &error);
Teraz czas na program w OpenCL. Program w C++ wykonuje sumowanie liczb sekwencyjnie. Najpierw sumuje pierwszą liczbę z wektora A z pierwszą liczbą z wektora B, następnie drugą liczbę z wektora A z drugą liczbą z wektora B itd.
Możemy łatwo zauważyć, że sumowanie pierwszych liczb nie jest w żaden sposób zależne od sumowania drugich liczb. Zamiast wykonywania tych operacji jedna po drugiej, możemy wykonać je równocześnie, co w żaden sposób nie wpłynie na wynik końcowy.
Poniższy program wykona sumowanie tylko dwóch liczb, a nie całego wektora. Skąd program wie, które liczby ma sumować? Funkcja „get_global_id” zwraca pewien unikalny identyfikator, na chwilę obecną możemy przyjąć, że to rozmiar wektora, który zdefiniowaliśmy w zmiennej „vectorSize”.
Umieszczamy poniższy kod w pliku „Add.cl”:
__kernel void Add(__global int* a, __global int* b, __global int* c, int size)
{
// Find position in global arrays.
int n = get_global_id(0);
// Bound check.
if (n < size)
{
c[n] = a[n] + b[n];
}
}
Wczytujemy kod programu z pliku, ładujemy go do GPU i kompilujemy:
// Read the OpenCL kernel in from source file.
std::ifstream file(".\\bin\\Add.cl", std::ifstream::in);
std::string str;
file.seekg(0, std::ios::end);
size_t programSize = (size_t)file.tellg();
str.reserve((unsigned int)file.tellg());
file.seekg(0, std::ios::beg);
str.assign(std::istreambuf_iterator<char>(file), std::istreambuf_iterator<char>());
const char* source = str.c_str();
// Create the program.
cl_program program = clCreateProgramWithSource(context, 1, &source, &programSize, &error);
error = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
Tworzymy kernel i przekazujemy do niego argumenty. Kernel to funkcja zdefiniowana w kodzie programu, który właśnie załadowaliśmy na GPU:
// Create the kernel.
cl_kernel kernel = clCreateKernel(program, "Add", &error);
// Set the Argument values.
error = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&bufferA);
error = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&bufferB);
error = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&bufferC);
error = clSetKernelArg(kernel, 3, sizeof(cl_int), (void*)&vectorSize);
Program jest gotowe, więc możemy przystąpić do wczytania wcześniej przygotowanych danych do GPU i uruchomienia kernela:
// Asynchronous write of data to GPU device.
error = clEnqueueWriteBuffer(commandQueue, bufferA, CL_FALSE, 0, sizeof(cl_int) * vectorSize, a, 0, NULL, NULL);
error = clEnqueueWriteBuffer(commandQueue, bufferB, CL_FALSE, 0, sizeof(cl_int) * vectorSize, b, 0, NULL, NULL);
// Launch kernel.
error = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, &vectorSize, &localWorkSize, 0, NULL, NULL);
// Read back results and check accumulated errors.
error = clEnqueueReadBuffer(commandQueue, bufferC, CL_TRUE, 0, sizeof(cl_int) * vectorSize, c, 0, NULL, NULL);
Funkcja „clEnqueueWriteBuffer” kopiuje dane z hosta (czyli naszego programu napisanego w C++) do GPU, a funkcja „clEnqueueReadBuffer” wykonuje operację odwrotną, czyli kopiuje dane z GPU do hosta.
Funckja „clEnqueueNDRangeKernel” uruchamia kernel. Jako piąty argument przekazujemy rozmiar wektóych, które sumujemy. Kernel zostanie uruchomiony tyle razy, ile wynosi wartość zmiennej „vectorSize”.
Teraz wracamy do kodu kernela i funkcji „get_global_id”. Jeżeli rozmiar wektora (vectorSize) to 32, to funkcja „get_global_id” w pierwszym uruchomionym kernelu zwróci 0, w drugim – 1, w trzecim – 2 itd. Dzięki temu jesteśmy w stanie określić które liczby z wektorów powinien sumować kernel.
Wyniki sumowania będą dostępne w zmiennej „c”. Zobaczmy czy operacja powiodła się:
for (size_t i = 0; i < vectorSize; ++i)
{
std::cout << a[i] << " + " << b[i] << " = " << c[i] << std::endl;
}
Pamiętamy, aby na koniec posprzątać po sobie i zwolnić pamięć:
// Cleanup and free memory.
clFlush(commandQueue);
clFinish(commandQueue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(bufferA);
clReleaseMemObject(bufferB);
clReleaseMemObject(bufferC);
clReleaseCommandQueue(commandQueue);
clReleaseContext(context);
delete[] a;
delete[] b;
delete[] c;
delete[] platformIds;
delete[] deviceIds;
Kod programu
Kod całego programu powinien wyglądać tak:
#include <iostream>
#include <fstream>
#include <string>
#include <CL/cl.h>
void randomizeArray(cl_int* data, size_t vectorSize)
{
for (size_t i = 0; i < vectorSize; ++i)
{
data[i] = rand() % 10;
}
}
int main()
{
cl_int error = CL_SUCCESS;
// Get platform number.
cl_uint platformNumber = 0;
error = clGetPlatformIDs(0, NULL, &platformNumber);
if (0 == platformNumber)
{
std::cout << "No OpenCL platforms found." << std::endl;
return 0;
}
// Get platform identifiers.
cl_platform_id* platformIds = new cl_platform_id[platformNumber];
error = clGetPlatformIDs(platformNumber, platformIds, NULL);
// Get platform info.
for (cl_uint i = 0; i < platformNumber; ++i)
{
char name[1024] = { '\0' };
std::cout << "Platform:\t" << i << std::endl;
error = clGetPlatformInfo(platformIds[i], CL_PLATFORM_NAME, 1024, &name, NULL);
std::cout << "Name:\t\t" << name << std::endl;
error = clGetPlatformInfo(platformIds[i], CL_PLATFORM_VENDOR, 1024, &name, NULL);
std::cout << "Vendor:\t\t" << name << std::endl;
std::cout << std::endl;
}
// Get device count.
cl_uint deviceNumber;
error = clGetDeviceIDs(platformIds[1], CL_DEVICE_TYPE_GPU, 0, NULL, &deviceNumber);
if (0 == deviceNumber)
{
std::cout << "No OpenCL devices found on platform " << 1 << "." << std::endl;
}
// Get device identifiers.
cl_device_id* deviceIds = new cl_device_id[deviceNumber];
error = clGetDeviceIDs(platformIds[1], CL_DEVICE_TYPE_GPU, deviceNumber, deviceIds, &deviceNumber);
// Get device info.
for (cl_uint i = 0; i < deviceNumber; ++i)
{
char name[1024] = { '\0' };
std::cout << "Device:\t\t" << i << std::endl;
error = clGetDeviceInfo(deviceIds[i], CL_DEVICE_NAME, 1024, &name, NULL);
std::cout << "Name:\t\t" << name << std::endl;
error = clGetDeviceInfo(deviceIds[i], CL_DEVICE_VENDOR, 1024, &name, NULL);
std::cout << "Vendor:\t\t" << name << std::endl;
error = clGetDeviceInfo(deviceIds[i], CL_DEVICE_VERSION, 1024, &name, NULL);
std::cout << "Version:\t" << name << std::endl;
}
std::cout << std::endl;
// Allocate and initialize host arrays
size_t vectorSize = 32;
size_t localWorkSize = 8;
cl_int* a = new cl_int[vectorSize];
cl_int* b = new cl_int[vectorSize];
cl_int* c = new cl_int[vectorSize];
randomizeArray(a, vectorSize);
randomizeArray(b, vectorSize);
// Create the OpenCL context.
cl_context context = clCreateContext(0, deviceNumber, deviceIds, NULL, NULL, NULL);
if (NULL == context)
{
std::cout << "Failed to create OpenCL context." << std::endl;
}
// Create a command-queue
cl_command_queue commandQueue = clCreateCommandQueue(context, deviceIds[0], 0, &error);
// Allocate the OpenCL buffer memory objects for source and result on the device.
cl_mem bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int) * vectorSize, NULL, &error);
cl_mem bufferB = clCreateBuffer(context, CL_MEM_READ_ONLY , sizeof(cl_int) * vectorSize, NULL, &error);
cl_mem bufferC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_int) * vectorSize, NULL, &error);
// Read the OpenCL kernel in from source file.
std::ifstream file(".\\bin\\Add.cl", std::ifstream::in);
std::string str;
file.seekg(0, std::ios::end);
size_t programSize = (size_t)file.tellg();
str.reserve((unsigned int)file.tellg());
file.seekg(0, std::ios::beg);
str.assign(std::istreambuf_iterator<char>(file), std::istreambuf_iterator<char>());
const char* source = str.c_str();
// Create the program.
cl_program program = clCreateProgramWithSource(context, 1, &source, &programSize, &error);
error = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
// Create the kernel.
cl_kernel kernel = clCreateKernel(program, "Add", &error);
// Set the Argument values.
error = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&bufferA);
error = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&bufferB);
error = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&bufferC);
error = clSetKernelArg(kernel, 3, sizeof(cl_int), (void*)&vectorSize);
// Asynchronous write of data to GPU device.
error = clEnqueueWriteBuffer(commandQueue, bufferA, CL_FALSE, 0, sizeof(cl_int) * vectorSize, a, 0, NULL, NULL);
error = clEnqueueWriteBuffer(commandQueue, bufferB, CL_FALSE, 0, sizeof(cl_int) * vectorSize, b, 0, NULL, NULL);
// Launch kernel.
error = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, &vectorSize, &localWorkSize, 0, NULL, NULL);
// Read back results and check accumulated errors.
error = clEnqueueReadBuffer(commandQueue, bufferC, CL_TRUE, 0, sizeof(cl_int) * vectorSize, c, 0, NULL, NULL);
// Print results.
for (size_t i = 0; i < vectorSize; ++i)
{
std::cout << a[i] << " + " << b[i] << " = " << c[i] << std::endl;
}
// Cleanup and free memory.
clFlush(commandQueue);
clFinish(commandQueue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseMemObject(bufferA);
clReleaseMemObject(bufferB);
clReleaseMemObject(bufferC);
clReleaseCommandQueue(commandQueue);
clReleaseContext(context);
delete[] a;
delete[] b;
delete[] c;
delete[] platformIds;
delete[] deviceIds;
// Press Enter, to quit application.
std::cin.get();
return 0;
}
Źródła w serwisie GitLab.
W następnej części kursu poznamy API dla języka C++.
