Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Task01 Василий Можаев SPbU #47

Closed
wants to merge 1 commit into from
Closed

Task01 Василий Можаев SPbU #47

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 7 additions & 2 deletions src/cl/aplusb.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@
// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять)
// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10)

__kernel void aplusb(...) {
__kernel void aplusb(__global const float as[], __global const float bs[], __global float cs[], unsigned int n) {
// Узнать, какой 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 All @@ -20,4 +20,9 @@ __kernel void aplusb(...) {
// и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС)
// поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы,
// либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой)
}

int id = get_global_id(0);
if (id >= n)
return;
cs[id] = as[id] + bs[id];
}
116 changes: 95 additions & 21 deletions src/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,19 +39,60 @@ int main() {

// TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь
// (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор)
int preferredDeviceType = 0;
cl_device_id preferredDevice;

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) {
auto device = devices[deviceIndex];
cl_device_type deviceType;
OCL_SAFE_CALL(clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(cl_device_type), &deviceType, nullptr));
if (deviceType == CL_DEVICE_TYPE_GPU) {
if (preferredDeviceType < 2) {
preferredDevice = device;
preferredDeviceType = 2;
}
} else if (deviceType == CL_DEVICE_TYPE_CPU) {
if (preferredDeviceType < 1) {
preferredDevice = device;
preferredDeviceType = 1;
}
}
}
}
if (preferredDeviceType == 0) {
throw std::runtime_error("Neither CPU nor GPU devices found!");
}

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

cl_int error;
auto context = clCreateContext(nullptr, 1, &preferredDevice, nullptr, nullptr, &error);
OCL_SAFE_CALL(error);

// Контекст и все остальные ресурсы следует освобождать с помощью clReleaseContext/clReleaseQueue/clReleaseMemObject... (да, не очень RAII, но это лишь пример)

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

unsigned int n = 1000 * 1000;
auto commandQueue = clCreateCommandQueue(context, preferredDevice, 0, &error);
OCL_SAFE_CALL(error);

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

auto bufsize = sizeof(float) * n;
auto abuf = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bufsize, as.data(), &error);
OCL_SAFE_CALL(error);
auto bbuf = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, bufsize, bs.data(), &error);
OCL_SAFE_CALL(error);
auto cbuf = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, bufsize, cs.data(), &error);
OCL_SAFE_CALL(error);

// TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl)
// затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания),
// напечатав исходники в консоль (if проверяет, что удалось считать хоть что-то)
Expand All @@ -79,36 +128,55 @@ 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* source_c_str = kernel_sources.c_str();
const std::size_t source_length = kernel_sources.size();
auto program = clCreateProgramWithSource(context, 1, &source_c_str, &source_length, &error);
OCL_SAFE_CALL(error);

// TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции
// см. clBuildProgram

error = clBuildProgram(program, 1, &preferredDevice, "", nullptr, nullptr);
bool success_build = true;
if (error != CL_BUILD_PROGRAM_FAILURE) OCL_SAFE_CALL(error);
else success_build = false;

// А также напечатайте лог компиляции (он будет очень полезен, если в кернеле есть синтаксические ошибки - т.е. когда 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, preferredDevice, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size));
std::vector<char> log(log_size, 0);
OCL_SAFE_CALL(clGetProgramBuildInfo(program, preferredDevice, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr));
if (log_size > 1) {
std::cout << "Log:" << std::endl;
std::cout << log.data() << std::endl;
}

if (!success_build)
return 0;

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

auto kernel = clCreateKernel(program, "aplusb", &error);
OCL_SAFE_CALL(error);

// 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*), &abuf);
clSetKernelArg(kernel, i++, sizeof(cl_mem*), &bbuf);
clSetKernelArg(kernel, i++, sizeof(cl_mem*), &cbuf);
clSetKernelArg(kernel, i++, sizeof(unsigned int), &n);
}

// TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности)
Expand All @@ -127,6 +195,9 @@ int main() {
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueNDRangeKernel...
// clWaitForEvents...
cl_event event;
OCL_SAFE_CALL(clEnqueueNDRangeKernel(commandQueue, kernel, 1, nullptr, &global_work_size, &workGroupSize, 0, nullptr, &event));
OCL_SAFE_CALL(clWaitForEvents(1, &event));
t.nextLap();// При вызове nextLap секундомер запоминает текущий замер (текущий круг) и начинает замерять время следующего круга
}
// Среднее время круга (вычисления кернела) на самом деле считается не по всем замерам, а лишь с 20%-перцентайля по 80%-перцентайль (как и стандартное отклонение)
Expand All @@ -140,34 +211,37 @@ int main() {
// - Флопс - это число операций с плавающей точкой в секунду
// - В гигафлопсе 10^9 флопсов
// - Среднее время выполнения кернела равно t.lapAvg() секунд
std::cout << "GFlops: " << 0 << std::endl;
std::cout << "GFlops: " << (double)n / t.lapAvg() / 1000000000 << 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: " << 3.0 * n * sizeof(float) / 1024 / 1024 / 1024 / t.lapAvg() << " GB/s" << std::endl;
}

// TODO 15 Скачайте результаты вычислений из видеопамяти (VRAM) в оперативную память (RAM) - из cs_gpu в cs (и рассчитайте скорость трансфера данных в гигабайтах в секунду)
{
timer t;
for (unsigned int i = 0; i < 20; ++i) {
// clEnqueueReadBuffer...
cl_event event;
OCL_SAFE_CALL(clEnqueueReadBuffer(commandQueue, cbuf, CL_TRUE, 0, bufsize, cs.data(), 0, nullptr, &event));
OCL_SAFE_CALL(clWaitForEvents(1, &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) / 1024 / 1024 / 1024 / t.lapAvg() << " 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;
}
Loading