Task02 Алсу Верещагина ITMO

ANSWER.md

@@ -0,0 +1,75 @@
+1) Пусть на вход дан сигнал x[n], а на выход нужно дать два сигнала y1[n] и y2[n]:
+
+ y1[n] = x[n - 1] + x[n] + x[n + 1]
+ y2[n] = y2[n - 2] + y2[n - 1] + x[n]
+Какой из двух сигналов будет проще и быстрее реализовать в модели массового параллелизма на GPU и почему?
+
+Ответ: Первый вариант проще и быстрее реализоать в модели массового параллелизма на GPU
+Объяснение: 
+- В варианте 1 сигнал y1 зависит от другого независимого сигнала x
+- В варианте 2 помимо независимого сигнала x присутсвует рекурсия, что требует обязательное вычисление предыдущих 
+значений
+
+
+2) Предположим что размер warp/wavefront равен 32 и рабочая группа делится на warp/wavefront-ы таким образом что 
+внутри warp/wavefront номер WorkItem по оси x меняется чаще всего, затем по оси y и затем по оси z.
+
+Напоминание: инструкция исполняется (пусть и отмаскированно) в каждом потоке warp/wavefront если хотя бы один поток 
+выполняет эту инструкцию неотмаскированно. Если не все потоки выполняют эту инструкцию неотмаскированно - происходит 
+т.н. code divergence.
+
+Пусть размер рабочей группы (32, 32, 1)
+
+int idx = get_local_id(1) + get_local_size(1) * get_local_id(0);
+if (idx % 32 < 16)
+ foo();
+else
+ bar();
+Произойдет ли code divergence? Почему?
+Ответ: code divergence не произойдет
+Объяснение: распишем переменную idx исходя из условия задачи:
+get_local_size(1) = 32
+idx = y + 32 * x;
+Следовательно выражение (idx % 32 < 16) будет зависеть лишь от y:
+(y + 32 * x) % 32 = y 
+При этом y внутри warp будет иметь константное значение
+
+
+3) Как и в прошлом задании предположим что размер warp/wavefront равен 32 и рабочая группа делится на 
+warp/wavefront-ы таким образом что внутри warp/wavefront номер WorkItem по оси x меняется чаще всего, затем по оси y 
+и затем по оси z.
+
+Пусть размер рабочей группы (32, 32, 1). Пусть data - указатель на массив float-данных в глобальной видеопамяти 
+идеально выравненный (выравнен по 128 байтам, т.е. data % 128 == 0). И пусть размер кеш линии - 128 байт.
+
+(a)
+
+data[get_local_id(0) + get_local_size(0) * get_local_id(1)] = 1.0f;
+Будет ли данное обращение к памяти coalesced? Сколько кеш линий записей произойдет в одной рабочей группе?
+
+
+Ответ: Обращение к памяти coalesed, 32 кэш линии
+Объяснение: data[x + 32 * y] = 1.0f; x принимает значения от 0 до 31.
+При этом 128 / 4 = 32 - кол-во данных типа float, которые помещаются в кэш.
+Следовательно для каждого значения y будет 1 кэш линия, для всего для рабочей группы будет 32 кэш линии
+
+
+(b)
+
+data[get_local_id(1) + get_local_size(1) * get_local_id(0)] = 1.0f;
+Будет ли данное обращение к памяти coalesced? Сколько кеш линий записей произойдет в одной рабочей группе?
+
+Ответ: Обращение к памяти не coalesed, 32 * 32 = 1024 кэш линии
+Объяснение: data[y + 32 * x] = 1.0f; x принимает значения от 0 до 31
+В данном случае обращение идет с шагом в 32. Кол-во записей кэш линий = 32 * 32 (количество warp) = 1024
+
+(c)
+
+data[1 + get_local_id(0) + get_local_size(0) * get_local_id(1)] = 1.0f;
+Будет ли данное обращение к памяти coalesced? Сколько кеш линий записей произойдет в одной рабочей группе?
+
+Ответ: Обращение к памяти coalesed, 33 кэш линии
+Объяснение: data[1 + x + 32 * y] = 1.0f; x принимает значения от 0 до 31
+Выражение аналогично тому, что в пункте а, однако в этом случае также идет сдвиг на +1, кол-во кэш линий = 32 (кол-во 
+warp) + 1 = 33
+