diff --git a/.figures/clion_edit_configurations.png b/.figures/clion_edit_configurations.png new file mode 100644 index 00000000..bd179b97 Binary files /dev/null and b/.figures/clion_edit_configurations.png differ diff --git a/.figures/clion_working_directory.png b/.figures/clion_working_directory.png new file mode 100644 index 00000000..8b8e732b Binary files /dev/null and b/.figures/clion_working_directory.png differ diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml index 4622bbbe..8af9ae05 100644 --- a/.github/workflows/cmake.yml +++ b/.github/workflows/cmake.yml @@ -7,7 +7,7 @@ env: jobs: build: - runs-on: ubuntu-18.04 + runs-on: ubuntu-20.04 steps: - uses: actions/checkout@v2 @@ -15,9 +15,9 @@ jobs: - name: Install prerequisites run: | sudo apt update - sudo apt install -y g++-5 - sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-5 90 - sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-5 90 + sudo apt install -y g++-7 + sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-7 90 + sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-7 90 - name: Install OpenCL driver for CPU run: sudo bash .github/scripts/install_intel_opencl.sh; @@ -30,4 +30,4 @@ jobs: - name: aplusb working-directory: ${{github.workspace}} - run: ./build/enumDevices \ No newline at end of file + run: ./build/aplusb diff --git a/CMakeLists.txt b/CMakeLists.txt index a4646666..1a0edce0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,9 +2,9 @@ cmake_minimum_required(VERSION 3.1) add_subdirectory(libs) -project(enumDevices) +project(aplusb) set(CMAKE_CXX_STANDARD 11) add_executable(${PROJECT_NAME} src/main.cpp) -target_link_libraries(${PROJECT_NAME} libclew) +target_link_libraries(${PROJECT_NAME} libclew libutils) diff --git a/README.md b/README.md index 9db0793e..2cb8ddc2 100644 --- a/README.md +++ b/README.md @@ -1,114 +1,21 @@ -В этом репозитории предложены задания для курса по вычислениям на видеокартах 2023. +В этом репозитории предложены задания для В этом репозитории предложены задания для курса по вычислениям на видеокартах 2023. [Остальные задания](https://github.com/GPGPUCourse/GPGPUTasks2023/). -# Задание 0. Вводное. +# Задание 1. A+B. -[![Build Status](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml/badge.svg?branch=task00&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml) - -Установка OpenCL-драйвера для процессора -======================================== - -Установить OpenCL-драйвер для процессора полезно, даже если у вас есть видеокарта, т.к. на нем удобно тестировать приложение (драйвер видеокарты гораздо чаще может повиснуть вместе с ОС). - -Windows -------- - -1. Откройте https://software.intel.com/content/www/us/en/develop/tools/opencl-cpu-runtime.html -2. Скачайте (требует регистрацию, [прямая ссылка для Windows](http://registrationcenter-download.intel.com/akdlm/irc_nas/vcp/13794/opencl_runtime_18.1_x64_setup.msi) - если не качает - попробуйте из-под инкогнито или [отсюда](https://disk.yandex.ru/d/dlVbMoI3tsPZfw)) -3. Установите - -Linux (Рекомендуется Ubuntu 18.04, 20.04 или 22.04) ----------------------------------- - -1. Откройте https://software.intel.com/content/www/us/en/develop/tools/opencl-cpu-runtime.html -2. Скачайте (требует регистрацию, [прямая ссылка для Ubuntu](http://registrationcenter-download.intel.com/akdlm/irc_nas/vcp/15532/l_opencl_p_18.1.0.015.tgz) - если не качает - попробуйте из-под инкогнито или [отсюда](https://disk.yandex.ru/d/dlVbMoI3tsPZfw)) -3. ``apt-get install -yq cpio`` -4. ``tar -xzf l_opencl_p_18.1.0.015.tgz`` -5. ``sudo ./l_opencl_p_18.1.0.015/install.sh`` -6. Проведите установку. - -Если у вас довольно новый процессор, например i7-8550U, то драйвер может его не поддерживать - ```clCreateContext``` вернет ошибку ```CL_DEVICE_NOT_AVAILABLE```, в таком случае поставьте свежий драйвер [отсюда](https://github.com/intel/compute-runtime/releases) (включает в т.ч. драйвер для встроенной Intel GPU). - -Если в процессе запуска этого задания процессор не виден как допустимое OpenCL-устройство - создайте **Issue** в этом репозитории с перечислением: - - - Версия OS - - Вывод команды ``ls /etc/OpenCL/vendors`` - - Если там в т.ч. есть ``intel.icd`` файл - то его содержимое (это маленький текстовый файл) - -Установка OpenCL-драйвера для видеокарты -======================================== - -Windows -------- - -Поставьте драйвер стандартным образом - скачав инсталлятор с официального сайта вендора вашей видеокарты и установив. - -Linux ------ - -NVidia: ``sudo apt install nvidia-<версия>`` (например, ``nvidia-384`` или ``nvidia-535``) - -AMD: [скачав](https://www.amd.com/en/support) и установив amdgpu-pro драйвер - -Проверка окружения и начало выполнения задания -============================================== - -Про работу под Windows см. в секции [Как работать под windows](#%D0%9A%D0%B0%D0%BA-%D1%80%D0%B0%D0%B1%D0%BE%D1%82%D0%B0%D1%82%D1%8C-%D0%BF%D0%BE%D0%B4-windows). - -1. Сделайте fork этого репозитория -2. ``git clone ВАШ_ФОРК_РЕПОЗИТОРИЯ`` -3. ``cd GPGPUTasks2023`` -4. ``git checkout task00`` -5. ``mkdir build`` -6. ``cd build`` -7. ``cmake ..`` -8. ``make -j4`` -9. ``./enumDevices`` должно увидеть хотя бы одну OpenCL-платформу: - -``` -Number of OpenCL platforms: 1 -Platform #1/1 - Platform name: -``` - -Если же вы видите ошибку: -``` -terminate called after throwing an instance of 'std::runtime_error' - what(): Can't init OpenCL driver! -Aborted (Core dumped) -``` -То попробуйте установить ```sudo apt install ocl-icd-libopencl1``` и выполнить ``./enumDevices`` снова. - -Если вы видите ошибку: -``` -: CommandLine Error: Option 'polly' registered more than once! -LLVM ERROR: inconsistency in registered CommandLine options -``` -То, наоборот, может помочь удалить пакет ```sudo apt remove ocl-icd-libopencl1``` и попробовать выполнить ``./enumDevices`` снова. - -Если ``./enumDevices`` не показывает хотя бы одну платформу - создайте **Issue** с перечислением: - - - OS, процессор и видеокарта - - Успешно ли прошла установка Intel-CPU драйвера - - Какое было поведение до установки пакета ``ocl-icd-libopencl1`` и какое поведение стало после - - Вывод ``./enumDevices`` +[![Build Status](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml/badge.svg?branch=task01&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2023/actions/workflows/cmake.yml) Задание ======= 0. Сделать fork проекта -1. Прочитать все комментарии подряд и выполнить все **TODO** в файле ``src/main.cpp``. Для разработки под Linux рекомендуется использовать CLion. Под Windows рекомендуется использовать CLion+MSVC. Также под Windows можно использовать Visual Studio Community. -2. Отправить **Pull-request** с названием ```Task00 <Имя> <Фамилия> <Аффиляция>```. **Аффиляция** - SPbU/HSE/ITMO. -3. В тексте **PR** укажите вывод программы при исполнении на сервере Github CI (Github Actions) и на вашем компьютере (в **pre**-тэгах, чтобы сохранить форматирование, см. [пример](https://raw.githubusercontent.com/GPGPUCourse/GPGPUTasks2023/task00/.github/pull_request_example.md)). И ваш бранч должен называться так же, как и у меня - **task00**. -4. Убедиться что Github CI (Github Actions) смог скомпилировать ваш код и что все хорошо (если нет - то поправить, пожалуйста, не используйте C++ из будущего, о котором не знает GCC 5.5) -5. Ждать комментарии проверки +1. Прочитать все комментарии подряд и выполнить все **TODO** в файле ``src/main.cpp`` и ``src/cl/aplusb.cl`` +2. Отправить **Pull-request** с названием```Task01 <Имя> <Фамилия> <Аффиляция>``` (добавив в описании вывод работы программы в **pre**-тэгах - см. [пример](https://raw.githubusercontent.com/GPGPUCourse/GPGPUTasks2023/task01/.github/pull_request_example.md)) -**Дедлайн**: 23:59 10 сентября. Но убедиться, что хотя бы одно OpenCL-устройство у вас обнаруживается, лучше как можно раньше, желательно, до начала лекции 8 сентября, чтобы было больше времени на решение проблем если они возникнут (см. **Проверка окружения** выше). +**Дедлайн**: 23:59 17 сентября. -Как работать под Windows -======================== +Коментарии +========== -1. Используйте **64-битный компилятор**, т.е. [amd64](/.figures/clion_msvc_settings.png), а не x86. (Если при запуске видите ``Invalid Parameter - 100``, то вы все еще используете 32-битный компилятор) -2. Рекомендуется использовать CLion+MSVC. -3. Можно использовать Visual Studio 2017 Community или новее, она поддерживает CMake-проекты (``File`` -> ``Open`` -> ``Cmake...``). Разве что передавать аргументы запускаемой программе [неудобно](https://docs.microsoft.com/en-us/cpp/ide/cmake-tools-for-visual-cpp?view=vs-2017#configure-cmake-debugging-sessions). +Т.к. в ``TODO 6`` исходники кернела считываются по относительному пути ``src/cl/aplusb.cl``, то нужно правильно настроить working directory. Например в случае CLion нужно открыть ``Edit configurations`` -> и указать ``Working directory: .../НАЗВАНИЕПАПКИПРОЕКТА`` (см. [подробнее](https://github.com/GPGPUCourse/GPGPUTasks2023/tree/task01/.figures)) diff --git a/libs/CMakeLists.txt b/libs/CMakeLists.txt index d50565b2..0b4ea7e9 100644 --- a/libs/CMakeLists.txt +++ b/libs/CMakeLists.txt @@ -1 +1,2 @@ add_subdirectory(clew) +add_subdirectory(utils) diff --git a/libs/utils/CMakeLists.txt b/libs/utils/CMakeLists.txt new file mode 100644 index 00000000..c6acc756 --- /dev/null +++ b/libs/utils/CMakeLists.txt @@ -0,0 +1,19 @@ +cmake_minimum_required(VERSION 3.1) + +project(libutils) + +set(HEADERS + libutils/fast_random.h + libutils/string_utils.h + libutils/timer.h + ) + +set(SOURCES + libutils/string_utils.cpp + ) + +set(CMAKE_CXX_STANDARD 11) + +add_library(${PROJECT_NAME} ${SOURCES} ${HEADERS}) +target_link_libraries(${PROJECT_NAME}) +target_include_directories(${PROJECT_NAME} PUBLIC ${PROJECT_SOURCE_DIR}) diff --git a/libs/utils/libutils/fast_random.h b/libs/utils/libutils/fast_random.h new file mode 100644 index 00000000..a03a077c --- /dev/null +++ b/libs/utils/libutils/fast_random.h @@ -0,0 +1,38 @@ +#pragma once + +#include + +// See https://stackoverflow.com/a/1640399 +class FastRandom { +public: + FastRandom(unsigned long seed=123456789) { + reset(seed); + } + + void reset(unsigned long seed=123456789) { + x = seed; + y = 362436069; + z = 521288629; + } + + // Returns pseudo-random value in range [min; max] (inclusive) + int next(int min=0, int max=std::numeric_limits::max()) { + x ^= x << 16; + x ^= x >> 5; + x ^= x << 1; + + unsigned long t = x; + x = y; + y = z; + z = t ^ x ^ y; + + return min + (unsigned int) (z % (((unsigned long) max) - min + 1)); + } + + float nextf() { + return (next() * 2000.0f / std::numeric_limits::max()) - 1000.0f; + } + +private: + unsigned long x, y, z; +}; diff --git a/libs/utils/libutils/string_utils.cpp b/libs/utils/libutils/string_utils.cpp new file mode 100644 index 00000000..ab828981 --- /dev/null +++ b/libs/utils/libutils/string_utils.cpp @@ -0,0 +1,158 @@ +#include "string_utils.h" +#include + +std::vector split(const std::string &string, const std::string &separator, bool keep_empty_parts) +{ + std::vector result; + size_t p = 0; + + while (true) { + size_t s = string.find(separator, p); + if (s == std::string::npos) + break; + std::string token = string.substr(p, s - p); + if (keep_empty_parts || token.size()) + result.push_back(token); + p = s + separator.size(); + } + + std::string token = string.substr(p); + if (keep_empty_parts || token.size()) + result.push_back(token); + return result; +} + +std::string join(const std::vector &tokens, const std::string &separator) +{ + std::string res; + for (size_t i = 0; i < tokens.size(); i++) { + if (i) + res += separator; + res += tokens[i]; + } + return res; +} + +std::istream &getline(std::istream &is, std::string &str) +{ + std::string::size_type nread = 0; + + if (std::istream::sentry(is, true)) { + std::streambuf *const sbuf = is.rdbuf(); + str.clear(); + + while (nread < str.max_size()) { + int c1 = sbuf->sbumpc(); + if (c1 == std::streambuf::traits_type::eof()) { + is.setstate(std::istream::eofbit); + break; + } else { + ++nread; + const char ch = c1; + if (ch != '\n' && ch != '\r') { + str.push_back(ch); + } else { + const char ch1 = is.peek(); + if (ch == '\n' && ch1 == '\r') is.ignore(1); + if (ch == '\r' && ch1 == '\n') is.ignore(1); + break; + } + } + } + } + + if (nread == 0 || nread >= str.max_size()) { + is.setstate(std::istream::failbit); + } + + return is; +} + +double atof(const std::string &s) +{ + std::stringstream ss(s); + ss.imbue(std::locale::classic()); + + double value = 0; + ss >> value; + return value; +} + +int atoi(const std::string &s) +{ + std::stringstream ss(s); + ss.imbue(std::locale::classic()); + + int value = 0; + ss >> value; + return value; +} + +std::string tolower(const std::string &str) +{ + std::string res = str; + for (size_t k = 0; k < res.size(); k++) res[k] = ::tolower(res[k]); + return res; +} + +std::string trimmed(const std::string &s) +{ + const size_t p1 = s.find_first_not_of(' '); + const size_t p2 = s.find_last_not_of(' '); + + if (p1 == std::string::npos) + return std::string(); + + return s.substr(p1, p2 - p1 + 1); +} + +// base 64 encoding/decoding +// http://stackoverflow.com/questions/180947/base64-decode-snippet-in-c + +std::string base64_encode(const std::string &in) +{ + std::string out; + + int val=0, valb=-6; + for (std::string::const_iterator it = in.begin(); it != in.end(); ++it) { + unsigned char c = *it; + + val = (val<<8) + c; + valb += 8; + while (valb>=0) { + out.push_back("ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"[(val>>valb)&0x3F]); + valb-=6; + } + } + + if (valb>-6) out.push_back("ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"[((val<<8)>>(valb+8))&0x3F]); + while (out.size()%4) out.push_back('='); + return out; +} + +std::string base64_decode(const std::string &in) +{ + std::string out; + + std::vector T(256,-1); + for (int i=0; i<64; i++) T["ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"[i]] = i; + + int val=0, valb=-8; + for (std::string::const_iterator it = in.begin(); it != in.end(); ++it) { + unsigned char c = *it; + if (isspace(c)) + continue; + + if (T[c] == -1) + break; + + val = (val<<6) + T[c]; + valb += 6; + if (valb>=0) { + out.push_back(char((val>>valb)&0xFF)); + valb-=8; + } + } + + return out; +} diff --git a/libs/utils/libutils/string_utils.h b/libs/utils/libutils/string_utils.h new file mode 100644 index 00000000..cbd96ec2 --- /dev/null +++ b/libs/utils/libutils/string_utils.h @@ -0,0 +1,24 @@ +#pragma once + +#include +#include +#include +#include + +template +std::string to_string(T value) +{ + std::ostringstream ss; + ss << value; + return ss.str(); +} + +std::vector split(const std::string &string, const std::string &separator, bool keep_empty_parts = true); +std::string join(const std::vector &tokens, const std::string &separator); +std::istream &getline(std::istream &is, std::string &str); +double atof(const std::string &s); +int atoi(const std::string &s); +std::string tolower(const std::string &str); +std::string trimmed(const std::string &str); +std::string base64_encode(const std::string &in); +std::string base64_decode(const std::string &in); diff --git a/libs/utils/libutils/timer.h b/libs/utils/libutils/timer.h new file mode 100644 index 00000000..fddf6d1a --- /dev/null +++ b/libs/utils/libutils/timer.h @@ -0,0 +1,161 @@ +#pragma once + +#ifdef _WIN32 +#include +#else +#include +#endif + +#include +#include +#include + +class timer { +protected: +#ifdef _WIN32 + typedef clock_t timer_type; +#else + typedef struct timeval timer_type; +#endif + + double counter_; + timer_type start_; + int is_running_; + + std::vector laps_; + +public: + timer(bool paused = false) + { + counter_ = 0; + is_running_ = 0; + if (!paused) + start(); + } + + void start() + { + if (is_running_) return; + + start_ = measure(); + is_running_ = 1; + } + + void stop() + { + if (!is_running_) return; + + counter_ += diff(start_, measure()); + is_running_ = 0; + } + + double nextLap() + { + double lap_time = elapsed(); + laps_.push_back(lap_time); + restart(); + return lap_time; + } + + void reset() + { + counter_ = 0; + is_running_ = 0; + } + + void restart() + { + reset(); + start(); + } + + double elapsed() const + { + double tm = counter_; + + if (is_running_) + tm += diff(start_, measure()); + + if (tm < 0) + tm = 0; + + return tm; + } + + const std::vector& laps() const + { + return laps_; + } + + // Note that this is not true averaging, if there is at least 5 laps - averaging made from 20% percentile to 80% percentile (See lapsFiltered) + double lapAvg() const + { + std::vector laps = lapsFiltered(); + + double sum = 0.0; + for (int i = 0; i < laps.size(); ++i) { + sum += laps[i]; + } + if (laps.size() > 0) { + sum /= laps.size(); + } + return sum; + } + + // Note that this is not true averaging, if there is at least 5 laps - averaging made from 20% percentile to 80% percentile (See lapsFiltered) + double lapStd() const + { + double avg = lapAvg(); + + std::vector laps = lapsFiltered(); + + double sum2 = 0.0; + for (int i = 0; i < laps.size(); ++i) { + sum2 += laps[i] * laps[i]; + } + if (laps.size() > 0) { + sum2 /= laps.size(); + } + return sqrt(std::max(0.0, sum2 - avg * avg)); + } + +protected: + + std::vector lapsFiltered() const + { + std::vector laps = laps_; + std::sort(laps.begin(), laps.end()); + + unsigned int nlaps = laps.size(); + if (nlaps >= 5) { + // Removing last 20% of measures + laps.erase(laps.end() - nlaps/5, laps.end()); + // Removing first 20% of measures + laps.erase(laps.begin(), laps.begin() + nlaps/5); + } + return laps; + } + + static timer_type measure() + { + timer_type tm; +#ifdef _WIN32 + tm = clock(); +#else + ::gettimeofday(&tm, 0); +#endif + return tm; + } + + static double diff(const timer_type &start, const timer_type &end) + { +#ifdef _WIN32 + return (double) (end - start) / (double) CLOCKS_PER_SEC; +#else + long secs = end.tv_sec - start.tv_sec; + long usecs = end.tv_usec - start.tv_usec; + + return (double) secs + (double) usecs / 1000000.0; +#endif + } +}; diff --git a/src/cl/aplusb.cl b/src/cl/aplusb.cl new file mode 100644 index 00000000..479624ac --- /dev/null +++ b/src/cl/aplusb.cl @@ -0,0 +1,23 @@ +#ifdef __CLION_IDE__ + // Этот include виден только для CLion парсера, это позволяет IDE "знать" ключевые слова вроде __kernel, __global + // а также уметь подсказывать OpenCL методы, описанные в данном инклюде (такие как get_global_id(...) и get_local_id(...)) + #include "clion_defines.cl" +#endif + +#line 8// Седьмая строчка теперь восьмая (при ошибках компиляции в логе компиляции будут указаны корректные строчки благодаря этой директиве) + +// TODO 5 реализуйте кернел: +// - От обычной функции кернел отличается модификатором __kernel и тем, что возвращаемый тип всегда void +// - На вход дано три массива float чисел; единственное, чем они отличаются от обычных указателей - модификатором __global, т.к. это глобальная память устройства (видеопамять) +// - Четвертым и последним аргументом должно быть передано количество элементов в каждом массиве (unsigned int, главное, чтобы тип был согласован с типом в соответствующем clSetKernelArg в T0D0 10) + +__kernel void aplusb(...) { + // Узнать, какой workItem выполняется в этом потоке поможет функция get_global_id + // см. в документации https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ + // OpenCL Compiler -> Built-in Functions -> Work-Item Functions + + // P.S. В общем случае количество элементов для сложения может быть некратно размеру WorkGroup, тогда размер рабочего пространства округлен вверх от числа элементов до кратности на размер WorkGroup + // и в таком случае, если сделать обращение к массиву просто по индексу=get_global_id(0), будет undefined behaviour (вплоть до повисания ОС) + // поэтому нужно либо дополнить массив данных длиной до кратности размеру рабочей группы, + // либо сделать return в кернеле до обращения к данным в тех WorkItems, где get_global_id(0) выходит за границы данных (явной проверкой) +} diff --git a/src/cl/clion_defines.cl b/src/cl/clion_defines.cl new file mode 100644 index 00000000..709ccae5 --- /dev/null +++ b/src/cl/clion_defines.cl @@ -0,0 +1,73 @@ +#ifndef clion_defines_cl // pragma once +#define clion_defines_cl + +#ifdef __CLION_IDE__ + +#define __kernel +#define __global +#define __local +#define __constant +#define __private + +#define half float + +struct float2 { float x; }; +struct float3 { float x, y, z; }; +struct float4 { float x, y, z, w; }; + +// https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/commonFunctions.html +#define gentype float +gentype clamp (gentype x, float minval, float maxval); +gentype degrees (gentype radians); +gentype max (gentype x, gentype y); +gentype min (gentype x, gentype y); +gentype mix (gentype x, gentype y, gentype a); +gentype radians (gentype degrees); +gentype sign (gentype x); +gentype smoothstep (gentype edge0, gentype edge1, gentype x); +gentype step (gentype edge, gentype x); +#undef gentype + +// https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/barrier.html +enum cl_mem_fence_flags +{ + CLK_LOCAL_MEM_FENCE, + CLK_GLOBAL_MEM_FENCE +}; +void barrier(cl_mem_fence_flags flags); + +// https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/vectorDataLoadandStoreFunctions.html +#define gentype float +#define gentypen float4 +gentypen vload4 (size_t offset, const gentype *p); +void vstore4 (gentypen data, size_t offset, gentype *p); +void vstore4 (gentypen data, size_t offset, gentype *p); +#undef gentypen +#undef gentype +float vload_half (size_t offset, const half *p); +float4 vload_half4 (size_t offset, const half *p); +void vstore_half (float data, size_t offset, half *p); +void vstore_half4 (float4 data, size_t offset, half *p); +float4 vloada_half4 (size_t offset, const half *p); +void vstorea_half4 (float4 data, size_t offset, half *p); + +// https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/workItemFunctions.html +uint get_work_dim (); +size_t get_global_size (uint dimindx); +size_t get_global_id (uint dimindx); +size_t get_local_size (uint dimindx); +size_t get_local_id (uint dimindx); +size_t get_num_groups (uint dimindx); +size_t get_group_id (uint dimindx); +size_t get_global_offset (uint dimindx); + +#ifndef STATIC_KEYWORD +#define STATIC_KEYWORD static +#endif + +#endif + +// 64 for AMD, 32 for NVidia, 8 for intel GPUs, 1 for CPU +#define WARP_SIZE 64 + +#endif // pragma once diff --git a/src/main.cpp b/src/main.cpp index e1acfd41..d70cde7b 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,6 +1,10 @@ #include #include +#include +#include +#include +#include #include #include #include @@ -29,64 +33,142 @@ void reportError(cl_int err, const std::string &filename, int line) { int main() { - // Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку libs/clew) + // Пытаемся слинковаться с символами OpenCL API в runtime (через библиотеку clew) if (!ocl_init()) throw std::runtime_error("Can't init OpenCL driver!"); - // Откройте - // https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ - // Нажмите слева: "OpenCL Runtime" -> "Query Platform Info" -> "clGetPlatformIDs" - // Прочитайте документацию clGetPlatformIDs и убедитесь, что этот способ узнать, сколько есть платформ, соответствует документации: - cl_uint platformsCount = 0; - OCL_SAFE_CALL(clGetPlatformIDs(0, nullptr, &platformsCount)); - std::cout << "Number of OpenCL platforms: " << platformsCount << std::endl; - - // Тот же метод используется для того, чтобы получить идентификаторы всех платформ - сверьтесь с документацией, что это сделано верно: - std::vector platforms(platformsCount); - OCL_SAFE_CALL(clGetPlatformIDs(platformsCount, platforms.data(), nullptr)); - - for (int platformIndex = 0; platformIndex < platformsCount; ++platformIndex) { - std::cout << "Platform #" << (platformIndex + 1) << "/" << platformsCount << std::endl; - cl_platform_id platform = platforms[platformIndex]; - - // Откройте документацию по "OpenCL Runtime" -> "Query Platform Info" -> "clGetPlatformInfo" - // Не забывайте проверять коды ошибок с помощью макроса OCL_SAFE_CALL - size_t platformNameSize = 0; - OCL_SAFE_CALL(clGetPlatformInfo(platform, CL_PLATFORM_NAME, 0, nullptr, &platformNameSize)); - // TODO 1.1 - // Попробуйте вместо CL_PLATFORM_NAME передать какое-нибудь случайное число - например 239 - // Т.к. это некорректный идентификатор параметра платформы - то метод вернет код ошибки - // Макрос OCL_SAFE_CALL заметит это, и кинет ошибку с кодом - // Откройте таблицу с кодами ошибок: - // libs/clew/CL/cl.h:103 - // P.S. Быстрый переход к файлу в CLion: Ctrl+Shift+N -> cl.h (или даже с номером строки: cl.h:103) -> Enter - // Найдите там нужный код ошибки и ее название - // Затем откройте документацию по clGetPlatformInfo и в секции Errors найдите ошибку, с которой столкнулись - // в документации подробно объясняется, какой ситуации соответствует данная ошибка, и это позволит, проверив код, понять, чем же вызвана данная ошибка (некорректным аргументом param_name) - // Обратите внимание, что в этом же libs/clew/CL/cl.h файле указаны всевоможные defines, такие как CL_DEVICE_TYPE_GPU и т.п. - - // TODO 1.2 - // Аналогично тому, как был запрошен список идентификаторов всех платформ - так и с названием платформы, теперь, когда известна длина названия - его можно запросить: - std::vector platformName(platformNameSize, 0); - // clGetPlatformInfo(...); - std::cout << " Platform name: " << platformName.data() << std::endl; - - // TODO 1.3 - // Запросите и напечатайте так же в консоль вендора данной платформы - - // TODO 2.1 - // Запросите число доступных устройств данной платформы (аналогично тому, как это было сделано для запроса числа доступных платформ - см. секцию "OpenCL Runtime" -> "Query Devices") - cl_uint devicesCount = 0; - - for (int deviceIndex = 0; deviceIndex < devicesCount; ++deviceIndex) { - // TODO 2.2 - // Запросите и напечатайте в консоль: - // - Название устройства - // - Тип устройства (видеокарта/процессор/что-то странное) - // - Размер памяти устройства в мегабайтах - // - Еще пару или более свойств устройства, которые вам покажутся наиболее интересными + // TODO 1 По аналогии с предыдущим заданием узнайте, какие есть устройства, и выберите из них какое-нибудь + // (если в списке устройств есть хоть одна видеокарта - выберите ее, если нету - выбирайте процессор) + + // TODO 2 Создайте контекст с выбранным устройством + // См. документацию https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/ -> OpenCL Runtime -> Contexts -> clCreateContext + // Не забывайте проверять все возвращаемые коды на успешность (обратите внимание, что в данном случае метод возвращает + // код по переданному аргументом errcode_ret указателю) + // И хорошо бы сразу добавить в конце clReleaseContext (да, не очень RAII, но это лишь пример) + + // 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; + // Создаем два массива псевдослучайных данных для сложения и массив для будущего хранения результата + std::vector as(n, 0); + std::vector bs(n, 0); + std::vector cs(n, 0); + FastRandom r(n); + for (unsigned int i = 0; i < n; ++i) { + as[i] = r.nextf(); + bs[i] = r.nextf(); + } + std::cout << "Data generated for n=" << n << "!" << std::endl; + + // TODO 4 Создайте три буфера в памяти устройства (в случае видеокарты - в видеопамяти - VRAM) - для двух суммируемых массивов as и bs (они read-only) и для массива с результатом cs (он write-only) + // См. Buffer Objects -> clCreateBuffer + // Размер в байтах соответственно можно вычислить через sizeof(float)=4 и тот факт, что чисел в каждом массиве n штук + // Данные в as и bs можно прогрузить этим же методом, скопировав данные из host_ptr=as.data() (и не забыв про битовый флаг, на это указывающий) + // или же через метод Buffer Objects -> clEnqueueWriteBuffer + // И хорошо бы сразу добавить в конце clReleaseMemObject (аналогично, все дальнейшие ресурсы вроде OpenCL под-программы, кернела и т.п. тоже нужно освобождать) + + // TODO 6 Выполните TODO 5 (реализуйте кернел в src/cl/aplusb.cl) + // затем убедитесь, что выходит загрузить его с диска (убедитесь что Working directory выставлена правильно - см. описание задания), + // напечатав исходники в консоль (if проверяет, что удалось считать хоть что-то) + std::string kernel_sources; + { + std::ifstream file("src/cl/aplusb.cl"); + kernel_sources = std::string(std::istreambuf_iterator(file), std::istreambuf_iterator()); + 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; + } + + // TODO 7 Создайте OpenCL-подпрограмму с исходниками кернела + // см. Runtime APIs -> Program Objects -> clCreateProgramWithSource + // у string есть метод c_str(), но обратите внимание, что передать вам нужно указатель на указатель + + // TODO 8 Теперь скомпилируйте программу и напечатайте в консоль лог компиляции + // см. clBuildProgram + + // А также напечатайте лог компиляции (он будет очень полезен, если в кернеле есть синтаксические ошибки - т.е. когда clBuildProgram вернет CL_BUILD_PROGRAM_FAILURE) + // Обратите внимание, что при компиляции на процессоре через Intel OpenCL драйвер - в логе указывается, какой ширины векторизацию получилось выполнить для кернела + // см. clGetProgramBuildInfo + // size_t log_size = 0; + // std::vector log(log_size, 0); + // 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 + + // 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++, ..., ...); + } + + // TODO 11 Выше увеличьте n с 1000*1000 до 100*1000*1000 (чтобы дальнейшие замеры были ближе к реальности) + + // TODO 12 Запустите выполнения кернела: + // - С одномерной рабочей группой размера 128 + // - В одномерном рабочем пространстве размера roundedUpN, где roundedUpN - наименьшее число, кратное 128 и при этом не меньшее n + // - см. clEnqueueNDRangeKernel + // - Обратите внимание, что, чтобы дождаться окончания вычислений (чтобы знать, когда можно смотреть результаты в cs_gpu) нужно: + // - Сохранить событие "кернел запущен" (см. аргумент "cl_event *event") + // - Дождаться завершения полунного события - см. в документации подходящий метод среди Event Objects + { + 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... + 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; + + // TODO 13 Рассчитайте достигнутые гигафлопcы: + // - Всего элементов в массивах по n штук + // - Всего выполняется операций: операция a+b выполняется n раз + // - Флопс - это число операций с плавающей точкой в секунду + // - В гигафлопсе 10^9 флопсов + // - Среднее время выполнения кернела равно t.lapAvg() секунд + std::cout << "GFlops: " << 0 << 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; } + // TODO 15 Скачайте результаты вычислений из видеопамяти (VRAM) в оперативную память (RAM) - из cs_gpu в cs (и рассчитайте скорость трансфера данных в гигабайтах в секунду) + { + timer t; + for (unsigned int i = 0; i < 20; ++i) { + // clEnqueueReadBuffer... + 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; + } + + // 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!"); + // } + // } + return 0; }