Task04 Степанов Николай SPbU

ANSWER.md

@@ -0,0 +1,7 @@
+# Транспонирование матрицы
+
+Я не получил особо интересных результатов. Наивная версия < С локальной памятью < С хорошим доступам к банкам. Интересно лишь то, что ускорение не особо настолько большое, как я ожидал. Примерно на 30% быстрее при избавлении от non-coalesced доступа (видимо тут решает кеш) и 1-2% при улучшении обращения к локальной памяти (зато это улучшение есть всегда, хотя оно не большое).
+
+# Умножение матриц
+
+Самый интересный результат здесь заключается в том, что последняя версия побеждает примерно в 2 раза при правильном подборе параметров (3.14 TFlops против 1.66 TFlops), при том, что количество инструкций уменьшилось точно меньше чем в 2 раза (каждые 8 load-ов и 4 fma превращаются в 5 load-ов и 4 fma). Получается, что есть что-то еще, что сильно влияет на ускорение. Скорее всего дело в кеше (изменился паттерн обращения к глобальной памяти => кеш стал лучше утилизироваться), но это просто теория на основе моих знаний о кеше на CPU. Я могу также предполагать, что дело в количестве переключений между рабочими группами, или в нагрузке на планировщик, или в ограничение скорости чтения из глобальной памяти.

src/cl/matrix_multiplication.cl

@@ -7,21 +7,109 @@
 
 // TILE_SIZE и WORK_PER_THREAD задаются через поле 'defines' в кернел конфиге
 
-__kernel void matrix_multiplication_naive()
-{
- // TODO
+__kernel void matrix_multiplication_naive(
+ __global float *as,
+ __global float *bs,
+ __global float *cs,
+ unsigned int m,
+ unsigned int k,
+ unsigned int n
+) {
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+
+ float sum = 0;
+ for (unsigned int l = 0; l < k; l++) {
+ sum += as[j * k + l] * bs[l * n + i];
+ }
+
+ cs[j * n + i] = sum;
 }
 
 #ifdef TILE_SIZE
-__kernel void matrix_multiplication_local()
-{
- // TODO
+__kernel void matrix_multiplication_local(
+ __global float *as,
+ __global float *bs,
+ __global float *cs,
+ unsigned int m,
+ unsigned int k,
+ unsigned int n
+) {
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+ unsigned int li = get_local_id(0);
+ unsigned int lj = get_local_id(1);
+
+ __local float tile_a[TILE_SIZE][TILE_SIZE];
+ __local float tile_b[TILE_SIZE][TILE_SIZE];
+
+ float sum = 0;
+ for (int base_k = 0; base_k < k; base_k += TILE_SIZE) {
+ tile_a[lj][li] = as[j * k + base_k + li];
+ tile_b[lj][li] = bs[(base_k + lj) * n + i];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ for (int l = 0; l < TILE_SIZE; l++) {
+ sum += tile_a[lj][l] * tile_b[l][li];
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ cs[j * n + i] = sum;
 }
 #endif
 
 #if defined(TILE_SIZE) && defined(WORK_PER_THREAD)
-__kernel void matrix_multiplication_local_wpt()
-{
- // TODO
+__kernel void matrix_multiplication_local_wpt(
+ __global float *as,
+ __global float *bs,
+ __global float *cs,
+ unsigned int m,
+ unsigned int k,
+ unsigned int n
+) {
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+ unsigned int li = get_local_id(0);
+ unsigned int lj = get_local_id(1);
+ unsigned int gi = get_group_id(0);
+ unsigned int gj = get_group_id(1);
+ const int WPT = WORK_PER_THREAD;
+ const int TPT = TILE_SIZE / WPT;
+
+ __local float tile_a[TILE_SIZE][TILE_SIZE];
+ __local float tile_b[TILE_SIZE][TILE_SIZE];
+
+ float sum[WPT];
+ for (int w = 0; w < WPT; w++) {
+ sum[w] = 0;
+ }
+
+ for (int base_k = 0; base_k < k; base_k += TILE_SIZE) {
+ for (int w = 0; w < WPT; w++) {
+ unsigned int wlj = w * TPT + lj;
+ unsigned int wj = wlj + gj * TILE_SIZE;
+ tile_a[wlj][li] = as[wj * k + base_k + li];
+ tile_b[wlj][li] = bs[(base_k + wlj) * n + i];
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ for (int l = 0; l < TILE_SIZE; l++) {
+ for (int w = 0; w < WPT; w++) {
+ unsigned int wlj = w * TPT + lj;
+ sum[w] += tile_a[wlj][l] * tile_b[l][li];
+ }
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ for (int w = 0; w < WPT; w++) {
+ unsigned int wlj = w * TPT + lj;
+ cs[(gj * TILE_SIZE + wlj) * n + i] = sum[w];
+ }
 }
 #endif

src/cl/matrix_transpose.cl

@@ -5,17 +5,47 @@
 
 #line 6
 
-__kernel void matrix_transpose_naive()
+#define TILE_SIZE 16
+
+__kernel void matrix_transpose_naive(__global float *a, __global float *a_t, unsigned int m, unsigned int k)
 {
- // TODO
+
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+ float x = a[j * k + i];
+ a_t[i * m + j] = x;
 }
 
-__kernel void matrix_transpose_local_bad_banks()
+__kernel void matrix_transpose_local_bad_banks(__global float *a, __global float *a_t, unsigned int m, unsigned int k)
 {
- // TODO
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+ unsigned int li = get_local_id(0);
+ unsigned int lj = get_local_id(1);
+ unsigned int gi = get_group_id(0);
+ unsigned int gj = get_group_id(1);
+
+ __local float tile[TILE_SIZE][TILE_SIZE];
+ tile[lj][li] = a[j * k + i];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ a_t[(gi * TILE_SIZE + lj) * m + gj * TILE_SIZE + li] = tile[li][lj];
 }
 
-__kernel void matrix_transpose_local_good_banks()
+__kernel void matrix_transpose_local_good_banks(__global float *a, __global float *a_t, unsigned int m, unsigned int k)
 {
- // TODO
+ unsigned int i = get_global_id(0);
+ unsigned int j = get_global_id(1);
+ unsigned int li = get_local_id(0);
+ unsigned int lj = get_local_id(1);
+ unsigned int gi = get_group_id(0);
+ unsigned int gj = get_group_id(1);
+
+ __local float tile[TILE_SIZE][TILE_SIZE + 1];
+ tile[lj][li] = a[j * k + i];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ a_t[(gi * TILE_SIZE + lj) * m + gj * TILE_SIZE + li] = tile[li][lj];
 }

src/main_matrix_multiplication.cpp

@@ -50,29 +50,26 @@ struct KernelConfig {
 
 KernelConfig makeNaiveConfig(unsigned int tile_size)
 {
- throw std::runtime_error("not implemented");
  std::string kernel_name = "matrix_multiplication_naive";
- gpu::WorkSize work_size(0, 0/*TODO*/);
+ gpu::WorkSize work_size(tile_size, tile_size, M, N);
  std::string defines;
  std::string prefix = "[naive, ts=" + std::to_string(tile_size) + "]";
  return KernelConfig{kernel_name, work_size, defines, prefix};
 }
 
 KernelConfig makeLocalConfig(unsigned int tile_size)
 {
- throw std::runtime_error("not implemented");
  std::string kernel_name = "matrix_multiplication_local";
- gpu::WorkSize work_size(0, 0/*TODO*/);
+ gpu::WorkSize work_size(tile_size, tile_size, M, N);
  std::string defines = "-DTILE_SIZE=" + std::to_string(tile_size);
  std::string prefix = "[local, ts=" + std::to_string(tile_size) + "]";
  return KernelConfig{kernel_name, work_size, defines, prefix};
 }
 
 KernelConfig makeLocalWPTConfig(unsigned int tile_size, unsigned int wpt)
 {
- throw std::runtime_error("not implemented");
  std::string kernel_name = "matrix_multiplication_local_wpt";
- gpu::WorkSize work_size(0, 0/*TODO*/);
+ gpu::WorkSize work_size(tile_size, tile_size / wpt, M, N / wpt);
  std::string defines = "-DTILE_SIZE=" + std::to_string(tile_size) + " -DWORK_PER_THREAD=" + std::to_string(wpt);
  std::string prefix = "[local wpt, ts=" + std::to_string(tile_size) + ", wpt=" + std::to_string(wpt) + "]";
  return KernelConfig{kernel_name, work_size, defines, prefix};
@@ -143,9 +140,6 @@ int main(int argc, char **argv)
 
  const std::vector<float> cs_cpu_reference = computeCPU(as.data(), bs.data());
 
- // TODO uncomment
- return 0;
-
  runTest(makeNaiveConfig(4), as.data(), bs.data(), cs_cpu_reference.data());
  runTest(makeNaiveConfig(8), as.data(), bs.data(), cs_cpu_reference.data());
  runTest(makeNaiveConfig(16), as.data(), bs.data(), cs_cpu_reference.data());

src/main_matrix_transpose.cpp

@@ -9,10 +9,12 @@
 #include <vector>
 #include <iostream>
 #include <stdexcept>
+#include <cassert>
 
 const int benchmarkingIters = 100;
 const unsigned int M = 4096;
 const unsigned int K = 4096;
+const unsigned int GROUP_SIZE = 16;
 
 void runTest(const std::string &kernel_name, const float *as)
 {
@@ -33,9 +35,10 @@ void runTest(const std::string &kernel_name, const float *as)
  // поставьте каретку редактирования кода внутри скобок конструктора WorkSize -> Ctrl+P -> заметьте что есть 2, 4 и 6 параметров
  // - для 1D, 2D и 3D рабочего пространства соответственно
 
- // TODO uncomment
-// gpu::WorkSize work_size(0, 0, 0, 0 /*TODO*/);
-// matrix_transpose_kernel.exec(work_size, as_gpu, as_t_gpu, M, K);
+ assert(M % GROUP_SIZE == 0);
+ assert(K % GROUP_SIZE == 0);
+ gpu::WorkSize work_size(GROUP_SIZE, GROUP_SIZE, K, M);
+ matrix_transpose_kernel.exec(work_size, as_gpu, as_t_gpu, M, K);
 
  t.nextLap();
  }
@@ -74,9 +77,6 @@ int main(int argc, char **argv)
  }
  std::cout << "Data generated for M=" << M << ", K=" << K << std::endl;
 
- // TODO uncomment
- return 0;
-
  runTest("matrix_transpose_naive", as.data());
  runTest("matrix_transpose_local_bad_banks", as.data());
  runTest("matrix_transpose_local_good_banks", as.data());