Skip to content

Commit

Permalink
task01 done
Browse files Browse the repository at this point in the history
  • Loading branch information
@vatican1
vatican1 committed Sep 14, 2023
1 parent 493f393 commit 12aee4b
Show file tree
Hide file tree
Showing 2 changed files with 108 additions and 28 deletions.
10 changes: 9 additions & 1 deletion src/cl/aplusb.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,15 @@
// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять)
// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10)

__kernel void aplusb(...) {
__kernel void aplusb(__global float *a,
__global float *b,
__global float *c,
unsigned int n)
{
int i = get_global_id(0);
if (i >= n)
return;
c[i] = a[i] + b[i];
// Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id
// см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/
// OpenCL Compiler -> Built-in Functions -> Work-Item Functions
Expand Down
126 changes: 99 additions & 27 deletions src/main.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,41 @@ void reportError(cl_int err, const std::string &filename, int line) {

#define OCL_SAFE_CALL(expr) reportError(expr, __FILE__, __LINE__)

std::pair<cl_platform_id, cl_device_id> getBetterDevice()
{
cl_uint platformsCount = 0;
OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platformsCount));
std::vector<cl_platform_id> platforms(platformsCount);
OCL_SAFE_CALL(clGetPlatformIDs(platformsCount, platforms.data(), nullptr));

for (int platformIndex = 0; platformIndex < platformsCount; ++platformIndex)
{
cl_platform_id platform = platforms[platformIndex];

cl_uint devicesCount = 0;
OCL_SAFE_CALL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &devicesCount));
std::vector<cl_device_id> devices(devicesCount);
OCL_SAFE_CALL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, devicesCount, devices.data(), nullptr));

for (int deviceIndex = 0; deviceIndex < devicesCount; ++deviceIndex)
{
cl_device_id device = devices[deviceIndex];
cl_device_type device_type;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(cl_device_type), &device_type, nullptr));
if (device_type & CL_DEVICE_TYPE_GPU)
{
std::cout << "find GPU!" << std::endl;
return std::make_pair(platform, device);
}
}
}
cl_uint devicesCount = 0;
OCL_SAFE_CALL(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, nullptr, &devicesCount));
std::vector<cl_device_id> devices(devicesCount);
OCL_SAFE_CALL(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, devicesCount, devices.data(), nullptr));
std::cout << "don't find GPU!" << std::endl;
return std::make_pair(platforms[0], devices[0]);
}

int main() {
// Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку clew)
Expand All @@ -40,18 +75,34 @@ int main() {
// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)

std::pair<cl_platform_id, cl_device_id> device = getBetterDevice();

cl_ulong deviceNameSize = 0;
OCL_SAFE_CALL(clGetDeviceInfo(device.second, CL_DEVICE_NAME, 0, nullptr, &deviceNameSize));
std::vector<unsigned char> deviceName(deviceNameSize, 0);
OCL_SAFE_CALL(clGetDeviceInfo(device.second, CL_DEVICE_NAME, deviceNameSize, deviceName.data(), nullptr));
std::cout << "DeviceName: " << deviceName.data() << std::endl;

// TODO 2 Создайте контекст с выбранным устройством
// См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Contexts -> clCreateContext
// Не забывайте проверять все возвращаемые коды на успешность (обратите внимание, что в данном случае метод возвращает
// код по переданному аргументом errcode_ret указателю)
// И хорошо бы сразу добавить в конце clReleaseContext (да, не очень RAII, но это лишь пример)

cl_int errcode_res = CL_SUCCESS;

cl_context context = clCreateContext(nullptr, 1, &device.second, nullptr, nullptr, &errcode_res);;
OCL_SAFE_CALL(errcode_res);

// TODO 3 Создайте очередь выполняемых команд в рамках выбранного контекста и устройства
// См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Runtime APIs -> Command Queues -> clCreateCommandQueue
// Убедитесь, что в соответствии с документацией вы создали in-order очередь задач
// И хорошо бы сразу добавить в конце clReleaseQueue (не забывайте освобождать ресурсы)

unsigned int n = 1000 * 1000;
cl_command_queue queue = clCreateCommandQueue(context, device.second, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &errcode_res);
OCL_SAFE_CALL(errcode_res);

unsigned int n = 1000 * 1000 * 100;
// Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата
std::vector<float> as(n, 0);
std::vector<float> bs(n, 0);
Expand All @@ -70,6 +121,14 @@ int main() {
// или же через метод Buffer Objects -> clEnqueueWriteBuffer
// И хорошо бы сразу добавить в конце clReleaseMemObject (аналогично, все дальнейшие ресурсы вроде OpenCL под-программы, кернела и т.п. тоже нужно освобождать)

cl_mem as_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, as.data(), &errcode_res);
OCL_SAFE_CALL(errcode_res);
cl_mem bs_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, bs.data(), &errcode_res);
OCL_SAFE_CALL(errcode_res);
cl_mem cs_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, cs.data(), &errcode_res);
OCL_SAFE_CALL(errcode_res);


// TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl)
// затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания),
// напечатав исходники в консоль (if проверяет, что удалось считать хоть что-то)
Expand All @@ -80,36 +139,45 @@ int main() {
if (kernel_sources.size() == 0) {
throw std::runtime_error("Empty source file! May be you forgot to configure working directory properly?");
}
// std::cout << kernel_sources << std::endl;
// std::cout << kernel_sources << std::endl;
}

// TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела
// см. Runtime APIs -> Program Objects -> clCreateProgramWithSource
// у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель
const char * kernel_sources_pointer = kernel_sources.c_str();
cl_program program = clCreateProgramWithSource(context, 1, &kernel_sources_pointer, nullptr, &errcode_res);
OCL_SAFE_CALL(errcode_res);

// TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции
// см. clBuildProgram
OCL_SAFE_CALL(clBuildProgram(program, 1, &device.second, nullptr, nullptr, nullptr));

// А также напечатайте лог компиляции (он будет очень полезен, если в кернеле есть синтаксические ошибки - т.е. когда clBuildProgram вернет CL_BUILD_PROGRAM_FAILURE)
// Обратите внимание, что при компиляции на процессоре через Intel OpenCL драйвер - в логе указывается, какой ширины векторизацию получилось выполнить для кернела
// см. clGetProgramBuildInfo
// size_t log_size = 0;
// std::vector<char> log(log_size, 0);
// if (log_size > 1) {
// std::cout << "Log:" << std::endl;
// std::cout << log.data() << std::endl;
// }
size_t log_size = 0;
OCL_SAFE_CALL(clGetProgramBuildInfo(program, device.second, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size));
std::vector<char> log(log_size, 0);
OCL_SAFE_CALL(clGetProgramBuildInfo(program, device.second, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr));
if (log_size > 1) {
std::cout << "Log:" << std::endl;
std::cout << log.data() << std::endl;
}

// TODO 9 Создайте OpenCL-kernel в созданной подпрограмме (в одной подпрограмме может быть несколько кернелов, но в данном случае кернел один)
// см. подходящую функцию в Runtime APIs -> Program Objects -> Kernel Objects

const char * kernel_name= "aplusb";
cl_kernel kernel = clCreateKernel(program, kernel_name, &errcode_res);
OCL_SAFE_CALL(errcode_res);
// TODO 10 Выставите все аргументы в кернеле через clSetKernelArg (as_gpu, bs_gpu, cs_gpu и число значений, убедитесь, что тип количества элементов такой же в кернеле)
{
// unsigned int i = 0;
// clSetKernelArg(kernel, i++, ..., ...);
// clSetKernelArg(kernel, i++, ..., ...);
// clSetKernelArg(kernel, i++, ..., ...);
// clSetKernelArg(kernel, i++, ..., ...);
unsigned int i = 0;
clSetKernelArg(kernel, i++, sizeof(cl_mem), &as_mem);
clSetKernelArg(kernel, i++, sizeof(cl_mem), &bs_mem);
clSetKernelArg(kernel, i++, sizeof(cl_mem), &cs_mem);
clSetKernelArg(kernel, i++, sizeof(unsigned int), &n);
}

// TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -125,50 +193,54 @@ int main() {
size_t workGroupSize = 128;
size_t global_work_size = (n + workGroupSize - 1) / workGroupSize * workGroupSize;
timer t;// Это вспомогательный секундомер, он замеряет время своего создания и позволяет усреднять время нескольких замеров
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueNDRangeKernel...
// clWaitForEvents...
for (unsigned int i = 0; i < 20; ++i)
{
cl_event start_kernel;
OCL_SAFE_CALL(clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, &start_kernel));
OCL_SAFE_CALL(clWaitForEvents(1, &start_kernel));
t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга
}
// Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение)
// подробнее об этом - см. timer.lapsFiltered
// P.S. чтобы в CLion быстро перейти к символу (функции/классу/много чему еще), достаточно нажать Ctrl+Shift+Alt+N -> lapsFiltered -> Enter
std::cout << "Kernel average time: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl;
std::cout << "Kernel average time: " << t.lapAvg() << " +- " << t.lapStd() << " s" << std::endl;

// TODO 13 Рассчитайте достигнутые гигафлопcы:
// - Всего элементов в массивах по n штук
// - Всего выполняется операций: операция a+b выполняется n раз
// - Флопс - это число операций с плавающей точкой в секунду
// - В гигафлопсе 10^9 флопсов
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "GFlops: " << 0 << std::endl;
std::cout << "GFlops: " << n / t.lapAvg() / 1e9<< std::endl;

// TODO 14 Рассчитайте используемую пропускную способность обращений к видеопамяти (в гигабайтах в секунду)
// - Всего элементов в массивах по n штук
// - Размер каждого элемента sizeof(float)=4 байта
// - Обращений к видеопамяти 2*n*sizeof(float) байт на чтение и 1*n*sizeof(float) байт на запись, т.е. итого 3*n*sizeof(float) байт
// - В гигабайте 1024*1024*1024 байт
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "VRAM bandwidth: " << 0 << " GB/s" << std::endl;
std::cout << "VRAM bandwidth: " << (double) (3 * n * sizeof(float)) / t.lapAvg() / (1024*1024*1024) << " GB/s" << std::endl;
}

// TODO 15 Скачайте результаты вычислений из видеопамяти (VRAM) в оперативную память (RAM) - из cs_gpu в cs (и рассчитайте скорость трансфера данных в гигабайтах в секунду)
{
timer t;
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueReadBuffer...
for (unsigned int i = 0; i < 20; ++i)
{
cl_event event;
OCL_SAFE_CALL(clEnqueueReadBuffer(queue, cs_mem, CL_TRUE, 0, n * sizeof(float), cs.data(), 0, nullptr, &event));
t.nextLap();
}
std::cout << "Result data transfer time: " << t.lapAvg() << "+-" << t.lapStd() << " s" << std::endl;
std::cout << "VRAM -> RAM bandwidth: " << 0 << " GB/s" << std::endl;
std::cout << "VRAM -> RAM bandwidth: " << (double)n * sizeof(float) / t.lapAvg() / (1024*1024*1024) << " GB/s" << std::endl;
}

// TODO 16 Сверьте результаты вычислений со сложением чисел на процессоре (и убедитесь, что если в кернеле сделать намеренную ошибку, то эта проверка поймает ошибку)
// for (unsigned int i = 0; i < n; ++i) {
// if (cs[i] != as[i] + bs[i]) {
// throw std::runtime_error("CPU and GPU results differ!");
// }
// }
for (unsigned int i = 0; i < n; ++i) {
if (cs[i] != as[i] + bs[i]) {
throw std::runtime_error("CPU and GPU results differ!");
}
}

return 0;
}

0 comments on commit 12aee4b

Please sign in to comment.