diff --git a/.images/modern-gpu-arch.png b/.images/modern-gpu-arch.png new file mode 100644 index 0000000..2958175 Binary files /dev/null and b/.images/modern-gpu-arch.png differ diff --git a/.images/pageable-pinnable-mem.png b/.images/pageable-pinnable-mem.png new file mode 100644 index 0000000..64bb893 Binary files /dev/null and b/.images/pageable-pinnable-mem.png differ diff --git a/.images/reduction-tree.png b/.images/reduction-tree.png new file mode 100644 index 0000000..44db3b8 Binary files /dev/null and b/.images/reduction-tree.png differ diff --git a/.images/sync-async-memcpy.png b/.images/sync-async-memcpy.png new file mode 100644 index 0000000..f26a2e9 Binary files /dev/null and b/.images/sync-async-memcpy.png differ diff --git a/.images/transaction-optimized-mem-access.png b/.images/transaction-optimized-mem-access.png new file mode 100644 index 0000000..1b824d8 Binary files /dev/null and b/.images/transaction-optimized-mem-access.png differ diff --git a/ANSWERS.md b/ANSWERS.md index 1fb3b9c..b550b75 100644 --- a/ANSWERS.md +++ b/ANSWERS.md @@ -1486,19 +1486,834 @@ for(int i = 0; i < 20; i++){ _Примечание: ответ формировался на основе наполнения лекции 04_. # 32. Задание ядра, потоки и блоки потоков на примере перемножения двух матриц в CUDA + +**Задание ядра** + +Для того, чтобы задать ядро необходимо написать `__global__`. Данное ключевое слово дает компилятору `nvcc` понять, +что данная функция предназначена для вызова из хоста для выполнения на девайсе. Эта функция и называется ядром. +Для задания дополнительных функций, вызываемых в ядре, необходимо создать функции с ключевым словом `__device__`. + +Пример задания ядра: + +```cuda +__global__ +void multiplyCuda(T* A, T* B, T* C, unsigned int N) +{ +... +} +``` + +**Задание потоков и блоков потоков** + +Видеокарта должна понимать сколько потоков и блоков потоков нужно использовать в ядре. +Для этого необходимо задать количество потоков и количество блоков потоков. Можно использовать +конкретные числа или задать с помощью типа `dim3`. Тип `dim3` - это целочисленный векторный тип, +основанный на `uint3`, который используется для указания размеров. +При определении переменной типа dim3 любой оставшийся неопределенным компонент инициализируется +значением 1. + +Пример создания `dim3`: + +```cuda +dim3 gridDim((N + TILE_SIZE - 1) / TILE_SIZE, (N + TILE_SIZE - 1) / TILE_SIZE); +dim3 blockDim(TILE_SIZE, TILE_SIZE); +``` + +В примере выше мы задаем размерность сетки(количество блоков для блочного умножения) и размерность блока. + +Пример вызова ядра из хоста: + +```cuda +multiplyCuda<<>>(d_a, d_b, d_c, N); +``` + +Полный код ядра: + +```cuda +__global__ +void multiplyCuda(T* A, T* B, T* C, unsigned int N) +{ + __shared__ float shared_A[TILE_SIZE][TILE_SIZE]; // блок в разделяемой памяти для блока из матрицы А + __shared__ float shared_B[TILE_SIZE][TILE_SIZE]; // блок в разделяемой памяти для блока из матрицы В + + // Вычисление глобального индекса в матрице + int globalRow = blockIdx.y * blockDim.y + threadIdx.y; + int globalCol = blockIdx.x * blockDim.x + threadIdx.x; + + float Cvalue = 0.0f; + + // Вычисление локального индекса в блоке + int row = threadIdx.y; + int col = threadIdx.x; + + // Умножение матрицы (блочный алгоритм) + for (int m = 0; m < (N + TILE_SIZE - 1) / TILE_SIZE; ++m) + { + // Если локальный индекс ряда внутри блока + if (row < N && (m * TILE_SIZE + col) < N) + { + shared_A[row][col] = A[globalRow * N + m * TILE_SIZE + col]; + } else + { + shared_A[row][col] = 0.0f; + } + // Если локальный индекс колонки внутри блока + if (col < N && (m * TILE_SIZE + row) < N) + { + shared_B[row][col] = B[(m * TILE_SIZE + row) * N + globalCol]; + } else + { + shared_B[row][col] = 0.0f; + } + __syncthreads(); + + // Вычисление умножения блоков + for (int k = 0; k < TILE_SIZE; ++k) + Cvalue += shared_A[row][k] * shared_B[k][col]; + + __syncthreads(); + } + + // Запись результата из разделяемой памяти в глобальную + if (globalRow < N && globalCol < N) + C[globalRow * N + globalCol] = Cvalue; +} +``` + # 33. Структура ядра и адресация на примере перемножения двух матриц в CUDA + +Структура ядра представлена в примере ниже. В данном случае для выполнения умножения +используется разделяемая память в которой происходит перемножение блоков из матрицы А и +матрицы В. + +После этого вычисляются глобальные индексы в матрице и локальные индексы в +вычисляемых блоках. Причем каждому потоку соответствует свой индекс, поскольку они вычисляются в +зависимости от позиции потока в блоке, размерности блока и позиции блока в сетке (`threadIdx.x|y|z`, `blockDim.x|y|z`, `blockIdx.x|y|z`). + +После этого блоки из A и B перемножаются между собой, а результат записывается +из разделяемой памяти в глобальную. + +```cuda +__global__ +void multiplyCuda(T* A, T* B, T* C, unsigned int N) +{ + __shared__ float shared_A[TILE_SIZE][TILE_SIZE]; // блок в разделяемой памяти для блока из матрицы А + __shared__ float shared_B[TILE_SIZE][TILE_SIZE]; // блок в разделяемой памяти для блока из матрицы В + + // Вычисление глобального индекса в матрице + int globalRow = blockIdx.y * blockDim.y + threadIdx.y; + int globalCol = blockIdx.x * blockDim.x + threadIdx.x; + + float Cvalue = 0.0f; + + // Вычисление локального индекса в блоке + int row = threadIdx.y; + int col = threadIdx.x; + + // Умножение матрицы (блочный алгоритм) + for (int m = 0; m < (N + TILE_SIZE - 1) / TILE_SIZE; ++m) + { + // Если локальный индекс ряда внутри блока + if (row < N && (m * TILE_SIZE + col) < N) + { + shared_A[row][col] = A[globalRow * N + m * TILE_SIZE + col]; + } else + { + shared_A[row][col] = 0.0f; + } + // Если локальный индекс колонки внутри блока + if (col < N && (m * TILE_SIZE + row) < N) + { + shared_B[row][col] = B[(m * TILE_SIZE + row) * N + globalCol]; + } else + { + shared_B[row][col] = 0.0f; + } + __syncthreads(); + + // Вычисление умножения блоков + for (int k = 0; k < TILE_SIZE; ++k) + Cvalue += shared_A[row][k] * shared_B[k][col]; + + __syncthreads(); + } + + // Запись результата из разделяемой памяти в глобальную + if (globalRow < N && globalCol < N) + C[globalRow * N + globalCol] = Cvalue; +} +``` + # 34. Синхронизация потоков, дивергенция потоков, функции голосования в CUDA. Примеры + +**Синхронизация потоков** + +Архитектура CUDA предлагает 2 основных способа синхронизации данных: + +- Синхронизация устройства, когда графический процессор гарантированно завершает поставленные на выполнение задачи, а работа центрального процессора для данного алгоритма блокируется до окончания работы GPU. При этом синхронизация может происходить как в явном виде, так и в неявном виде +- Синхронизация нитей внутри блока. + +В качестве синхронизации нитей внутри блока CUDA может предложить: +- `void __syncthreads()`. +- `int __syncthreads_count(int predicate)`. Число нитей, для которых выражение predicate не равно 0. +- `int __syncthreads_and(int predicate)`. Не нулевое значение только в том случае, если значение predicate для всех нитей блока не равно 0. +- `int __syncthreads_or(int predicate)`. Не нулевое значение только в том случае, если значение predicate хотя бы для одной нити блока не равно 0. + +Варп представляет собой блок из 32 нитей мультипроцессора GPU, которые выполняют +одну и ту же инструкцию, но каждый над своим блоком данных, т.е. фактически варп +определяет минимальное число нитей, которые будут работать в один момент времени. +Это свойство дает ему возможность сделать инструкцию `__syncthreads()` лишней в коде. + +Пример синхронизации в ядре: + +```cuda +__global__ +void kernel(T* a, T* b, T* c, int N) +{ + for (int i = 0; i < N; i++) + { + // Что то выполняется потоками... + __syncthreads(); + } +} +``` + +Пример синхронизации вне ядра: +```cuda +kernel<<>>(d_a, d_b, d_c, N); +cudaDeviceSynchronize(); +``` + +**Дивергенция потоков** + +Дивергенция потоков относится к ситуации, когда потоки в варпе отклоняются +от одного и того же потока управления, влияя на параллелизм. + +Возникает, когда в тредах одного варпа инструкция условного перехода отрабатывает +по-разному, то есть в части тредов переход выполняется, а в остальной части нет. + +Пример принудительной дивергенции: + +```cuda +__global__ +void kernel(T* a, T* b, T* c, int N) +{ + for (int i = 0; i < threadIdx.x; i++) + { + // Каждый поток из варпа выйдет из цикла на разных итерациях цикла. + } +} +``` + +Время выполнения такого варпа будет равно суммарным временем исполнения всех выполненных ветвей. + +**Функции голосования** + +Начиная с CUDA 2.х доступны функции «голосования» между потоками варпа: +- `__all(bool) = 1` у всех потоков варпа bool=1 +- `__any(bool) = 1` у одного из потоков варпа bool=1 +- `__ballot(bool) = int`, побитово собраны значения bool потоков + +Начиная с CUDA 3.х доступны функции «перестановки» +значений, расположенных на регистрах потоков варпа: +- `__shfl(int var, int src) = значению var из потока src` +var – своё значение, src – номер потока, из которого читать + # 35. Архитектура современного GPU + +Современный GPU состоит из нескольких потоковых мультипроцессоров (SMs), каждый из которых состоит из +нескольких ядер с общим управлением и памятью. + +![alt text](.images/modern-gpu-arch.png) + +Потокам/блокам требуются ресурсы для выполнения (например, регистры, память, и т.д.), +поэтому SM может обслуживать ограниченное количество потоков/блоков одновременно. +Оставшиеся блоки ожидают завершения работы других блоков, +прежде чем их можно будет назначить SM. + +Потоки, находящиеся в разных блоках не могут, например: +- Не синхронизируются по `__syncthreads();` +- Не могут использовать одну и ту же разделяемую память + +Потоки из одного блока назначены строго одному потоковому мультипроцессору. +Потоки назначаются SMам одновременно. Блок не может быть назначен SM до тех пор, +пока он не обеспечит достаточное количество ресурсов для выполнения всех своих потоков. + +Потоки, назначенные SM, выполняются одновременно: +- В SM есть планировщик, который управляет их выполнением + +Блоки, назначенные SM, далее делятся на варпы, которые являются единицей планирования. +Размер варпов зависит от устройства, но на сегодняшний день он всегда 32 потоков/варп. + +Потоки в warp планируются совместно и выполняются в соответствии с моделью SIMD: +- Одна команда, несколько данных +- Одна команда извлекается и выполняется всеми потоками в warp, + каждый из которых обрабатывает разные данные + +Виды памяти, использующиеся в GPU: +- Глобальная +- Разделяемая +- Регистровая +- Текстурная +- Локальная + # 36. Понятие occupancy в CUDA. Пример расчета + +Потоковые мультипроцессоры (SM) — это система, состоящая из нескольких планировщиков варпов. + +Occupancy SM - это отношение активных потоков варпа к максимально возможному. + +Пример расчета на архитектуре Volta V100 (1024 потоков/блок, 2048 потоков/SM, 32 блоков/SM): + +Если 256 потоков/блок активно, то occupancy максимальный. +- (2048 потока)/(256 потоков/блок) = 8 блоков < 32 блока и 32 % 8 = 0 + +Если 32 потоков/блок активно, то occupancy будет ограничен из-за ограничения допустимого количества блоков к мультипроцессору. +- (2048 потока)/(32 потоков/блок) = 64 блоков < 32 блока + +Если 768 потоков/блок активно, то occupancy будет ограничен из-за того, что 2048 на делиится на 768 нацело. +И оставшиеся 512 потоков не могут быть использованы. +- (2048 потока)/(768 потоков/блок) = 2 блоков < 32 блока (1536 задействовано, 512 потоков простаивает). + # 37. Типы памяти в CUDA. Примеры создания и организации доступа + +Технология CUDA использует следующие типы памяти: регистры, локальная, +глобальная, разделяемая, константная и текстурная память. + +Разработчику недоступны к управлению регистры и локальная память. +Все остальные типы он вправе использовать по своему усмотрению. + +**Глобальная память** + +Глобальная память – тип памяти с самой высокой латентностью, из доступных на GPU. + +Переменные в данном типе памяти можно выделить с помощью +спецификатора `__global__`, а так же динамически, с помощью функций из семейства +`cudaMalloc`. + +Глобальная память в основном служит для хранения больших объемов данных, над которыми осуществляется обработка, +и для сохранения результата работы. + +Пример создания: + +```cuda +// Матрица А на девайсе +T* d_a = nullptr; + +// Выделение памяти на девайсе +cudaMalloc(&d_a, M * N * sizeof(T)); +``` + +Пример доступа: + +```cuda +__global__ +void kernel(T* a, T* b, T* c, int N) +{ + ... + a[row * N + col] = 1; + b[row * N + col] = 2; + c[row * N + col] = 3; + ... +} +``` + +**Разделяемая память** + +Разделяемая память относится к типу памяти с низкой латентностью. + +Данный тип памяти рекомендуется использовать для минимизации обращения к +глобальной памяти. Адресация разделяемой памяти осуществляется между нитями +одного блока, что может быть использовано для обмена данными между потоками +в пределах одного блока. + +Для размещения данных в разделяемой памяти используется спецификатор `__shared__`. + +Пример создания и доступа: + +```cuda +__global__ +void kernel(T* a, T* b, T* c, int N) +{ + __shared__ float smem[TILE_SIZE][TILE_SIZE]; + ... + // Преодоление проблемы с доступом к банкам памяти + smem[threadIdx.x] = a[blockIdx.x * 512 * 3 + threadIdx.x]; + smem[threadIdx.x + 512] = a[blockIdx.x * 512 * 3 + threadIdx.x + 512]; + smem[threadIdx.x + 1024] = a[blockIdx.x * 512 * 3 + threadIdx.x + 1024]; + ... +} +``` + +**Константная память** + +Константная память является одной из самых быстрых из доступных на GPU. + +Отличительной особенностью данного типа памяти является возможность записи данных +с хоста, но при этом в пределах GPU возможно лишь чтение из этой памяти, +что и обуславливает её название. + +Для размещения данных в константной памяти предусмотрен спецификатор `__constant__`. + +Для записи с хоста в константную память используется функция `cudaMemcpyToSymbol`. + +Пример создания и доступа: + +```cuda +__constant__ T a; // в глобальной области (не в мейне) + +cudaMemcpyToSymbol(a, b, N * M * sizeof(T), 0); // в мейне + +for (int i = 0; i < a; i++) // в ядре +``` + +**Текстурная память** + +Текстурная память входит в состав текстурных блоков, используемых в графических +задачах для формирования текстур. + +В текстурном блоке аппаратно реализована фильтрация текстурных координат, +интерполяция, нормализация текстурных координат в случаях, +когда они выходят за допустимые пределы. + +Текстурная память выделяется с помощью функции `cudaMallocArray` и +освобождается с использованием `cudaFreeArray`: + +```cuda +cudaError_t cudaMallocArray ( cudaArray_t* array, const cudaChannelFormatDesc* desc, size_t width, size_t height = 0, unsigned int flags = 0 ); + +cudaError_t cudaFreeArray ( cudaArray_t array ); +``` + +Для того чтобы получить доступ к выделенной памяти в CUDA-ядре требуется +ассоциировать указатель с текстурной ссылкой: + +```cuda +cudaError_t cudaBindTexture ( size_t* offset, const textureReference* texref, const void* devPtr, const cudaChannelFormatDesc* desc, size_t size = UINT_MAX ); +``` + +**Регистровая память** + +Регистровая память является самой быстрой из всех типов памяти. +Определить общее количество регистров, доступных GPU, можно с помощью +функции `cudaGetDeviceProperties`. + +**Локальная память** + +Локальная память может быть использована компилятором, +если все локальные переменные не могут быть размещены в регистровой памяти. +По скоростным характеристикам локальная память значительно медленнее, чем регистровая. + # 38. Механизм транзакций в CUDA. Пример + +Поскольку самая минимальная транзакция - 4 байта, то наиболее эффективным способом +доступа к памяти будет являться доступ одного потока к последовательным четырем байтам памяти. + +![alt text](.images/transaction-optimized-mem-access.png) + +Пример кода: + +```cuda +__global__ +void filterGPU(unsigned* inp, unsigned* outp, int pitch, int channels) +{ + __shared__ unsigned i_smem[w * w]; + __shared__ unsigned char buff_smem[w * w * 4]; + unsigned char* uc_smem = reinterpret_cast(i_smem); + ... + // Запись в разделяемую память интами + i_smem[destY * w + destX] = inp[(baseY * pitch) / 4 + baseX]; + ... + // Работаем с unsigned char* + buff_smem[base1D] = uc_smem[multY * (threadIdx.y + i) + x1D + j * channels + c]; + ... + // Возвращаем результат работы + baseY = blockIdx.y * TILE_WIDTH + threadIdx.y; + baseX = blockIdx.x * TILE_WIDTH + threadIdx.x; + outp[(baseY * pitch) / 4 + baseX] = reinterpret_cast(buff_smem)[threadIdx.x + threadIdx.y * w]; +} +``` + # 39. Конфликт по банкам в разделяемой памяти в CUDA. Пример + +Разделяемая память делится на блоки фиксированного размера, доступ к +которым осуществляется одновременно. + +Таким образом, если нити обращаются к различным банкам, +латентность доступа будет минимальна. Если хотя бы одна из нитей не получила +данные из-за конфликта доступа к банку, осуществляется +повторный доступ к памяти и данные догружаются. + +Данная процедура будет повторяться до тех пор, пока не будут разрешены +все конфликты и загружены или сохранены все данные. +В результате время доступа будет произведением +порядка конфликта на латентность доступа к разделяемой памяти. + +Пример: + +```cuda +__global__ +void kernel(T* a, T* b, int N) +{ + __shared__ T smem[512 * 3]; + // копируем входные данные в разделяемую память + smem[threadIdx.x] = a[blockIdx.x * 512 * 3 + threadIdx.x]; + smem[threadIdx.x + 512] = a[blockIdx.x * 512 * 3 + threadIdx.x + 512]; + smem[threadIdx.x + 1024] = a[blockIdx.x * 512 * 3 + threadIdx.x + 1024]; + __syncthreads(); + ... + // находим сумму и сохраняем результат работы + b[blockIdx.x * 512 + threadIdx.x] = smem[threadIdx.x * 3] + + smem[threadIdx.x * 3 + 1] + + smem[threadIdx.x * 3 + 2]; +} +``` + # 40. Алгоритм редукции в CUDA. Пример + +Редукция это процесс преобразования множества значений в одно результирующее значение. +Примером может являться сумма/произведение/разность и т.д. всех элементов массива. + +Сложность реализации алгоритма на CUDA является необходимость построения дерева редукции +на разделяемой памяти. + +![alt text](.images/reduction-tree.png) + +Пример редукции на CUDA: + +```cuda +__global__ +void reduceBase(int* inputData, int* outputData) +{ + extern __shared__ int smemData[]; + // загружаем данные в разделяемую память + const unsigned tid = threadIdx.x; + smemData[tid] = inputData[blockIdx.x * blockDim.x+ threadIdx.x]; + // синхронизируемся + __syncthreads(); + // выполняем редукцию над элементами массива + for (unsigned s = 1; s < blockDim.x; s <<= 1) + { + if(tid % (s << 1) == 0) + { + smemData[tid] += smemData[tid + s]; + } + // синхронизируемся, чтобы гарантировать + // корректность входных данных + // на следующей итерации цикла + __syncthreads(); + } + // сохраняем результат вычислений + if (tid == 0) outputData[blockIdx.x] = smemData[0]; +} +``` + # 41. Алгоритм свертки в CUDA. Пример + +Свертка - операция над парой матриц A (размера Nx×Ny) и B (размера Mx×My), +результатом которой является матрица C=A∗B размера (Nx−Mx+1)×(Ny−My+1). + +Каждый элемент результата вычисляется как скалярное произведение матрицы B +и некоторой подматрицы A такого же размера. + +Пример свертки на CUDA: + +```cuda +__constant__ +char gkernel[9] = +{ + 1, 0, -1, + 1, 0, -1, + 1, 0, -1 +}; + +__global__ +void convGPU(unsigned char* inp, unsigned char* outp, int width, int height) +{ + __shared__ unsigned char N_ds[w][w]; + int dest = threadIdx.y * TILE_WIDTH + threadIdx.x, + destY = dest / w, destX = dest % w, + srcY = blockIdx.y * TILE_WIDTH + destY, + srcX = blockIdx.x * TILE_WIDTH + destX, + src = (srcY * width + srcX) * channels + c; + if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width) + N_ds[destY][destX] = inp[src]; + else + N_ds[destY][destX] = 0; + + dest = threadIdx.y * TILE_WIDTH + threadIdx.x + TILE_WIDTH * TILE_WIDTH; + destY = dest / w, destX = dest % w; + srcY = blockIdx.y * TILE_WIDTH + destY; + srcX = blockIdx.x * TILE_WIDTH + destX; + src = (srcY * width + srcX) * channels + c; + if (destY < w) + { + if (srcY >= 1 && srcY < height && srcX >= 1 && srcX < width) + N_ds[destY][destX] = inp[src]; + else + N_ds[destY][destX] = 0; + } + __syncthreads(); + + int buff = 0; + int y, x; + for (y = 0; y < 3; y++) + { + for (x = 0; x < 3; x++) + { + buff += N_ds[threadIdx.y + y][threadIdx.x + x] * gkernel[y * 3 + x]; + } + } + y = blockIdx.y * TILE_WIDTH + threadIdx.y; + x = blockIdx.x * TILE_WIDTH + threadIdx.x; + if (y < height && x < width) + { + outp[y * width + x] = buff < 0 ? 0 : buff > 255 ? 0 : buff; + } +} +``` + # 42. Алгоритм операции инклюзивного scan в CUDA. Пример + +Инклюзивный скан - операция сложения всех элементов массива предшествующих некоторому +элементу j включительно. + +Пример: +[0, 1, 2, 3, 4, 5] -> [0, 1, 3, 6, 10, 15] + +```cuda +__global__ +void inc_scan(T* input, T* output, unsigned long long len) +{ + unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; + __shared__ float buffer_s[BLOCK_DIM]; + buffer_s[threadIdx.x] = input[i]; + __syncthreads(); + for (unsigned int stride = 1; stride <= BLOCK_DIM/2; stride *= 2) + { + float v; + if (threadIdx.x >= stride) + { + v = buffer_s[threadIdx.x - stride]; + } + __syncthreads(); + if (threadIdx.x >= stride) + { + buffer_s[threadIdx.x] += v; + } + __syncthreads(); + } + if(threadIdx.x == BLOCK_DIM - 1) + { + partialSums[blockIdx.x] = buffer_s[threadIdx.x]; + } + output[i] = buffer_s[threadIdx.x]; +} +``` + # 43. Алгоритм операции эксклюзивного scan в CUDA. Пример + + +Эксклюзивный скан - операция сложения всех элементов массива предшествующих некоторому +элементу j, не включающий сам элемент j. + +Пример: +[0, 1, 2, 3, 4, 5] -> [0, 0, 1, 3, 6, 10] + +```cuda +__global__ +void exc_scan(T* input, T* output, unsigned long long len) +{ + unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; + __shared__ float buffer1_s[BLOCK_DIM]; + __shared__ float buffer2_s[BLOCK_DIM]; + float* inBuffer_s = buffer1_s; + float* outBuffer_s = buffer2_s; + if (threadIdx.x == 0) + { + inBuffer_s[threadIdx.x] = 0.0f; + } + else + { + inBuffer_s[threadIdx.x] = input[i - 1]; + } + __syncthreads(); + for (unsigned int stride = 1; stride <= BLOCK_DIM/2; stride *= 2) + { + if(threadIdx.x >= stride) + { + outBuffer_s[threadIdx.x] = inBuffer_s[threadIdx.x] + inBuffer_s[threadIdx.x - stride]; + } + else + { + outBuffer_s[threadIdx.x] = inBuffer_s[threadIdx.x]; + } + __syncthreads(); + float* tmp = inBuffer_s; + inBuffer_s = outBuffer_s; + outBuffer_s = tmp; + } + if(threadIdx.x == BLOCK_DIM - 1) + { + partialSums[blockIdx.x] = inBuffer_s[threadIdx.x] + input[i]; + } + output[i] = inBuffer_s[threadIdx.x]; +} +``` + # 44. Асинхронное и синхронное копирование в CUDA. Pinned память. Способы выделения + +Время выполнения задачи = время работы ядра + обмен данными между CPU и +GPU. + +Для копирования памяти без участия ЦП применяется zero-copy операции. При +этом неявно используется копирование памяти сначала в pinned буфер, потом в +DRAM GPU. Разработчик может самостоятельно организовать логику получения +копируемой на GPU pinned памяти, что позволяет ускорить процесс копирования. +Это проявляется при вызове асинхронных функций копирования памяти. +Основная разница во времени выполнения асинхронных и синхронных версий +команд копирования проявляется при работе с pinned памятью в качестве входного +параметра. + +Pinned-память - page-locked память, т.е. страницы виртуальной памяти становятся +жёстко привязаны к физической. С такой памятью возможны асинхронные +операции с перемещением памяти host->dev/dev->host. + +Копирование такой памяти происходит быстрее, т.к. при копировании pageable +памяти под капотом происходит перенос её содержимого в pinned буфер. С такой +памятью работает DMA, из-за чего и достигается мгновенное возвращение из +функций асинхронного копирования. Её возможно не только выделять, но и делать +таковой уже выделенную pageable память. + +![alt text](.images/pageable-pinnable-mem.png) + +![alt text](.images/sync-async-memcpy.png) + +Для выделения памяти в pinned-режиме достаточно выделить память в ОЗУ с применением следующей функции: + +```cuda +cudaError_t cudaMallocHost( + void** ptr, // указатель на выделяемую память в ОЗУ + size_t size // объем выделяемой памяти +); +``` + +Заранее выделенная системная память может быть также зарегистрирована +в pinned-режиме с использованием CUDA API: + +```cuda +cudaError_t cudaHostRegister( + void * ptr, // указатель на память ОЗУ + size_t size, // размер регистрируемой памяти в байтах + unsigned int flags +); +``` + # 45. CUDA Stream. Создание, инициализация и синхронизация + +Технология `CUDA Stream` позволяет при соблюдении ряда условий (как +аппаратных, так и програмных) организовать параллельное выполнение некоторых +задач. Особенно удобно при работе с более чем одним gpu. + +Все команды, принимающие опциональный параметр типа `cudaStream` +по-умолчанию используют поток по умолчанию. +Команды из разных потоков, отличных от потока по-умолчанию, могут исполняться +параллельно в зависимости от аппаратных возможностей. + +Возможные случаи: +- Параллельные копирование и выполнение ядра +- Параллельные выполнение ядер +- Параллельные копирования с хоста на устройство и с устройства на хост + +Возможности параллельного выполнения можно проверить при получении свойств +устройства (структура `cudaDeviceProp`) + +- Если `cudaDeviceProp::asyncEngineCount > 0` устройство может выполнять +параллельно копирование и счет ядра + - Хостовая память должна быть page-locked + +```cuda +cudaMallocHost(&aHost , size); +cudaStreamCreate(&stream1); +cudaStreamCreate(&stream2); +cudaMemcpyAsync(aDev, aHost, size, cudaMemcpyHostToDevice, stream1); +kernel <<>>(...); +``` + +- Если `cudaDeviceProp::concurrentKernels > 0` устройство может выполнять ядра +параллельно + +```cuda +cudaStreamCreate(&stream1); +cudaStreamCreate(&stream2); +kernel1 <<>>(data_1); +kernel2 <<>>(data_2); +``` + +- Если `cudaDeviceProp::asyncEngineCount == 2` устройство может выполнять +параллельно копирование в обе стороны и счет ядра + +```cuda +cudaMallocHost(&aHost, size); +cudaMallocHost(&bHost, size); +cudaMemcpyAsync(aDev, aHost, size, cudaMemcpyHostToDevice, stream1); +cudaMemcpyAsync(bHost, bDev, size, cudaMemcpyDeviceToHost, stream2); +kernel <<>>(...); +``` + +**Неявная синхронизация** + +- Неявная синхронизация (ожидание завершения всех команд на устройтве) +выполняется перед: + - Выделением page-locked памяти / памяти на устройстве + - cudaMemSet + - Копированием между пересекающимися областями памяти на устройстве + - Отправкой команды в поток по-умолчанию + - Переключением режима кеша L1 +- Если между отправкой двух команд в разные потоки стоит что-то из этого списка параллельного выполнения не будет + +**Синхронизация по событию** + +- `cudaError_t cudaEventQuery(cudaEvent_t event)` + - Возвращает `cudaSuccess`, если событие уже произошло (вся работа до последнего + `cudaEventRecord` выполнена): иначе `cudaErrorNotReady` +- `cudaError_t cudaEventSynchronize (cudaEvent_t event)` + - Возвращает управление хостовой нити только после наступления события + +**Синхронизация по GPU** + +- cudaError_t cudaStreamWaitEvent +- (cudaStream_t stream, cudaEvent_t event, unsigned int flags ) +- Команды, отправленные в stream, начнут выполняться после наступления +события event + - Синхронизация будет эффективно выполнена на GPU + - При stream == NULL будут отложены все команды всех потоков +- Событие event может быть записано на другом GPU + - Синхронизация между GPU + +Пример: + +```cuda +A1<<<1, 1, 0, streamA>>>(d); +cudaEventRecord(halfA, streamA); +cudaStreamWaitEvent(streamB, halfA, 0); +B1<<<1, 1, 0, streamB>>>(d); +cudaEventRecord(halfB, streamB); +cudaStreamWaitEvent(streamA, halfB, 0); +A2<<<1, 1, 0, streamA>>>(d); +B2<<<1, 1, 0, streamB>>>(d); +``` + +**Синхронизация по потоку** + +- `cudaError_t cudaStreamQuery (cudaStream_t stream);` + - Возвращает `cudaSuccess`, если выполнены все команды в потоке stream, иначе + `cudaErrorNotReady` + +- `cudaError_t cudaStreamSynchronize (cudaStream_t stream);` + - Возвращает управление хостовой нити, когда завершится выполнение всех + команд, отправленных в поток stream + # 46. Микроархитектура Intel Knights Landing и ее наследники # 47. Микроархитектура Intel Knights Mill # 48. Микроархитектура Intel Sunny Cove