added task01

.github/workflows/cmake.yml

@@ -30,4 +30,4 @@ jobs:
 
     - name: aplusb
       working-directory: ${{github.workspace}}
-      run: ./build/enumDevices
+      run: ./build/aplusb

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)

README.md

@@ -2,113 +2,20 @@
 
 [Остальные задания](https://github.com/GPGPUCourse/GPGPUTasks2024/).
 
-# Задание 0. Вводное.
+# Задание 1. A+B.
 
-[![Build Status](https://github.com/GPGPUCourse/GPGPUTasks2024/actions/workflows/cmake.yml/badge.svg?branch=task00&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2024/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 lsb-core``
-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 GPGPUTasks2024``
-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/GPGPUTasks2024/actions/workflows/cmake.yml/badge.svg?branch=task01&event=push)](https://github.com/GPGPUCourse/GPGPUTasks2024/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/GPGPUTasks2024/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/GPGPUTasks2024/task01/.github/pull_request_example.md))
 
-**Дедлайн**: 23:59 10 сентября. Но убедиться, что хотя бы одно OpenCL-устройство у вас обнаруживается, лучше как можно раньше, желательно, до начала лекции 8 сентября, чтобы было больше времени на решение проблем если они возникнут (см. **Проверка окружения** выше).
+**Дедлайн**: 23:59 22 сентября.
 
-Как работать под 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/GPGPUTasks2024/tree/task01/.figures))

libs/CMakeLists.txt

@@ -1 +1,2 @@
 add_subdirectory(clew)
+add_subdirectory(utils)

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) выходит за границы данных (явной проверкой)
+}

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