Cuda

Warp - пачка потоков которые физически выполняются параллельно (32-е штуки). Управления warp-ми выполняет gpu.

Каждый поток в рамках warp-а имеет свой собственный ALU, но все они делят общий Control Unit. 

threadId - у потоков в рамках warp образуют непрерывный отрезок индексов. Во время выполнения warp физически находится на чипе и с него сходит только после окончания выполнения. Поэтому переключения между варпами ничего не стоит.

В терминах некоторых специальный инструкций имеется так называемый LaneId по факту это threadId.x % 32 для одномерной топологии блоков.

Selected Warp - варп который сейчас исполняется

Eligible Warp - варп готовый к исполнению. В каждый момент времени, чтобы SM не простаивал должен быть хотя бы один готовый к запуску warp. Этим прячется латентность вычислений

Stalled Warp - варп который не доступен для исполнения из-за неготовности входных данных

SM - Stream Multiprocessor. Процессор построенный по идеологии SIMD, но с наворотами:

 1. В каждый момент времени только один warp выполняется SM-ом.

 2. В warp следующая инструкция которого имеет все входы становится легальной для планирования.

 3. Все потоки в пределах варпа выполняют одну инструкцию

 4. Если нет потоков для выполнения - то GPU ничего не выполняет. Наиболее популярная причина - ожидание доступа в   

      глобальную память. NVIDIA вместо SIMD использует термин SIMT так как в принципе возможно (хоть это не очень 

      эффективно). Исполнения разных инструкций потоками в пределах warp.

"Производительность - то же самое что параллелизм, эффективность это тоже самое что локальность. С энергетической точки зрения большая энергия приходящая на чип тратится на передвижение данных" - Bill Daily

Про архитектуру GPU и выполнение CUDA кода на ней

GPU содержит:

  - Свою память (несколько гигабайт)

  - Ряд потоковых и независимых друг от друга потоковых мультипроцессоров (SM) (в районе 30 штук)

Память на GPU

Register File - все регистры потокового мультипроцессора. Обычные автоматические переменные обычно хранятся как регистры. Массив тоже может храниться в регистровом файле, если доступ осуществляется по известным на этапе компиляции индексам. В Fermi доступно 63 регистров на поток, в Kepler расширено на 255 регистров на поток. [3, стр.138]

Local Memory - все автоматические переменные которые не помещаются в регистры засовываются сюда. На самом деле это просто глобальная память, несмотря на название.

Shared Memory - находится на чипе. Нечто похожее на CPU L1 cache. Захватывается при исполнении ThreadBlock-а, освобождается по окончанию выполнения ThreadBlock-а. cudaFuncSetCacheConfig

Константная - 64Кб ведёт наилучшим образом когда все потоки варпа читают с одного адреса, память кешируется.

Текстурная - как и константная, но есть доп.аппартные штуки для работы с текстурой - фильтрование и т.д.

Глобальная - память гпу, актуальна во время всей работы приложения. Запись в одно и тоже место из разных потоков ведёт к undefined behaviour.

Кеши - L1 на SM, L2 на все SM, Read-only текстурная и константная. На GPU кеши работают только на load. На store не работают.

Zero-copy - pined виртуальная память процесса доступная на устройстве. Очень медленная скорость. Однако может пригодится при нехватке памяти. Можно ожидать относительно дешёвого копирования в интегрированных GPU находяшимся на одном чипе с CPU. При наличии отдельных (дискретных GPU) преимущества от zero-copy имеется только в специальных случаях.

Типы памяти

Для GPU кешируется только доступ в текcтурную и константную память.

На хосте можно выделить page-locked или pinned память через cudaMallocHost(...). 

При асинхронных операциях копирования CPU<>GPU использовать надо именно её (но это не точно).

Все обращение мультипроцессора к памяти происходит отдельно для каждого half-warp-а. 

Латентность (Cycles) / Throughput (Operations/Cycle) / Латентность  X Throughput (Little Law)

  ALU for Fermi - 20 CYCLES / 32 OP per CYCLE / 640 OP

  ALU for Kepler- 20 CYCLES / 192 OP per CYCLE / 3 840 OP

  MEM.OP for Fermi - 800 CYCLES / 144Gb per SEC or 92 Bytes per CYCLE / 74KB

  MEM.OP for Kepler- 800 CYCLES / 250Gb per SEC or 96 Bytes per CYCLE / 77KB

Типовые ограничения на GPU, каждая характеристика даёт оценку сверху

[3, p.89] [3, [.77]

  ThreadsPerBlock                 -- 1024

  ConcurentBlocks PerSM    -- 8/16

  ConcurentWarps PerSM    -- 48/64

  ConcurentThreads PerSM -- 1536/2048

  32-bit registers PerSM -- 32K/64K

  32-bit registers perThread-- 63/255

  SharedMemPerSM -- 48K

  Int and float ALU PerSM -- 32/48/192

  SpecialFunctUnit for float PerSM -- 4/8/32

  WarpScheulersPerSM -- 2/4

  Load/StoreUnitsPerSM -- 16/32

  L2 cache -- 768K/1536K

  GlobalMemoryPerDevice -- 6GB/12GB

Обычный патерн работы с CUDA

1. Разделяем аккуратно данные на блоки. 

RO (read-only) данные отправляем(заливаем) в __constant__ память.

 Лучше всего разделять даже не входные данные, а выходные. Это совет от людей проектирующих GPU из [2].

2. Выделяем память на GPU, копируем туда входные данные. CPU=>GPU

3. Запускаем ядро.

3.1. Если нужен RW доступ к общим данным в пределах блока то используем быструю __shared__ память. Для синхронизации работы по RW с shared памятью надо использовать __syncthreads. Если не все потоки придут в эту инструкцию в пределах блока то это в худшем случае приведёт к зависанию всей машины. 

Фишка CUDA не обеспечивать никакие примитивы синхронизации между блоками. В пределах одного варпа у всех потоков есть неявная синхронизация на уровне аппаратуры. Все доступы потоков в глобальную или shared память будут видны после этой точки синхронизации для потоков после этой точки.

3.2 Лучше использоваться барьеры при работе с shared памятью, если возможно. Если шаблон доступа к этой памяти потоков в пределах блока "разный" то возможно целесообразнее использовать atomic*. 

3.2. Если нужен RW доступ к общим данным в пределах потока используем регистры. Увы если нужны индексы в пределах блока то придётся использовать медленную локальную память. К регистрам принципиально не возможно обращаться по индексам изменяющимся во время выполнения.

3.3. atomic* на shared память гораздо быстрее чем в глобальную память. Сериалазацию потоков на доступ к переменной выполняет GPU.

4. Копируем выходные данные с GPU=>CPU через cudaMemCpy. Очищаем память на GPU

cudaMemCpy для копирования использует DMA контроллер. Если получиться лучше делать так что в каждый момент времени - 

* Cpu было чем-то занято

* DMA был занят копированием

* GPU был занят выполнением какого-либо ядра вашей программы.

Про неправильные вещи

1. __shared__ int* ptr; -- указатель хранится в shared памяти, это не указатель на shared память. Указатели в CUDA не содержат привязка к пространству памяти.

2. Лучше избегать указателей на указатели. Архитектура GPU не расчитана на самом деле на конструкции типа связанных списков и т.д. Лучше при работе с указателями отдавать предпочтение простым паттернам доступа.

Про характеристики вычислительной программы

[1] Количество load - ов на поток. (Байты)

[2] Кол-во float операций на поток. (Штук операций)

[3] GMAC - Global Memory Access to Flop Ratio = [1]/[2]. (байт/операцию)

[4] Необходимое ширина канала доступа в память, чтобы достичь пиковой производительности - GMAC  * PeakFlops

      PeakFlops для GTX 260 - 805GFlops.

[5] Реальная ширина канала(Actual BandWidth). Например для GTX 260 112 GB/sec

[6] Какая реальная производительность в GFlops для приложения - "Actual BandWidth/GMAC "

Про SM-ы на GPU

SM цель:

-- Вообще цель создание программы для gpu - пускать вычислительные потоки на gpu. 

-- Архитектурное решение всей платформы - потоки группируются в блоки

SM Ограничения:

-- Может выполнять одновременно 1024 потока (32 варпа по 32 потока каждый)

-- Может выполнять до 8 блоков одновременно. Для загрузки GPU нужно иметь 128-192 threads на ThreadBlock

-- Блок целиком выполняется на одном мультипроцессоре (Stream Multiprocessor)

-- Разбиение потоков на warp-ы происходит независимо для каждого блока

-- Потоки могут взаимодействовать через барьерную синхронизацию и разделяемую память только с потоками в пределах одного блока

SM содержит:

-- Special Function Unit, Instruction Unit, Double Precision Unit 

-- Память (Register File и Shared Memory). Она делится между всеми резидентными ThreadBlocks на GPU

-- Кеш константой и текстурной памяти доступной на чтение Scalar Processor-ом

-- Несколько скалярных ядер (Scalar Processor). Каждая нить выполняется на одном из скалярном процессоров.

Все варпы физически выполняются независимо, поэтому если разные warp-ы пошли по разным if/else, то ничего страшного. 

Для одного варпа - плохо (divergence, branching). 

Будут выполнены все ветки, с аппаратной пометной кто выполняется, а какой из потоков варпа простаивает на данный момент.

-- Так же в [9] отмечено что для больших веток будет вставлено так называемое warp voting, чтобы вообще не тратить время

на декодирование команд команд которые не выполняются вообще не одним потоком.

-- Так же в [9] отмечено, что проверки часто связаны с границей области на которой происходит расчёт. В этом случае рекомендация такая:

подумать на сколько сложны вычисления для различия границы и внутренности области задачи.

Возможно даже стоить написать два ядра - одно для обработки внутренней области, а другое для обработки границы.

-- Так же в [8] отмечено, что вообще говоря если обработка элементов включает обработку "сложных" и "простых" ситуаций одновременно то лучше разделить логику.

Общее для всех SM:

  -- (Tex). Для доступа к текстурной памяти используется текстурный блок, используемый несколькими мультипроцессорами

  -- (Tex + Sm -ы его используюшие) образуют так называемый TPC (Texture Processing Cluster)

  -- GPU с точки зрения CUDA представляет собой набор TPC (Tpc Array)

  -- Zero-overhad на запуск ядра, на переключение контекста потока

  -- Барьеры дешёвые, не следует использовать спин-блокировку

Про строительные блоки параллельной разработки

scan(aka prefix sum) - используется в radix sort, run-length-encoding, histogram, string comparision, others.

scan(a)=[0,a0+a1,a0+a1+a2,... ]

Параллельная реализация для блока shared памяти - каждый поток в пределах блока ассоциируется с элементом в префиксной сумме, поток вычисляет значение в своей ячейке со значением в ячейке со смещением threadId.x - o. потоки синхронизируются барьером. 

Производят обновление/запись shared памяти, и снова синхронизируется. o умножается на два. Всего такой цикл повториться log2(N) раз. Алгоритм работает inplace нет необходимости во втором буфере.

segemented scan - scan + barrier flags to "restart" scan - сделать много scans за один вызов. реализация похожа на обычный scan.

Итерационные методы решения СЛАУ проще ложаться на GPU чем прямые методы. Пара примеров итерационных методов:

-- Гауса-Зейделя (используется новые элементы для решения каждого уравнения). Сходится для симметрично положительно-определенной матрицы.

-- Якоби (для вычисления каждого любого нового элемента используются только элементы предыдущей итерации). Сходится для симметрично положительно-определенной матрицы с диаг. преобладанием относительно строк.

Оптимизация, общие подходы

  1. Максимизация параллелизма

  2. Оптимизация доступа в память

  3. Оптимизация математики

Оптимизация, сложные приёмы

В VS при использовании nvcc можно выставить --ptxas-options = -v. 

В output VS будет выводиться информация о кол-ве регистров, используемой локальной, разделяемой и константной памяти.

При использовании nvcc для того чтобы это увидеть можно использовать nvcc -Xptxas  -v. После этого полученные числа можно засунуть в CUDA Occupancy Calculator

Чтобы ограничить использование регистров можно использовать следующий флаг nvcc --maxrregcount=x

Лучше запускать на выполнение независимые ядра.

Чтобы произошло объединения запроса на чтение памяти (coalescing) для всего half-warp-а для его нитей требуется выполнение след. условий:

   1. Все 32/2 нитей обращаются к 4-байтным(32 бита) словам, давая в результате 64 байтный блок. Или все 32/2 нитей       

      обращаются к 8-байтным(64 бита) словам, давая в резльтате 128 байтный блок.

  2. Блок выровнен по своему размеру. Адрес 64-ех байтного блока кратен 64, 128-байтного блока кратен 128

  3. Все 16 слов лежат в пределах блока

  4. Нити обращаются к словам последовательно. (Для CC 1.2 уже не обязательно)

Как правило лучше использовать StructureOfArrays, а не ArrayOfStructures.

Лучше снабжать указатели модификаторами const, restrict для убирания алиасинга.

Компиляция и разработка программ на CUDA

1. Расширения файлов *.cu и они компилируются через nvcc, у которого много ключей

2. CUDA host api -- апи только для CPU кода и может использоваться в зависимости от задачи в двух режимах

  2.1 CUDA driver api (низкоуровневый) все имена функций есть cu*

  2.2 CUDA runtime api (высокоуровневый) все имена функций есть cuda*

Флажки компилятора

-XCompiler -- специфицирует опции которые передаются компилятору под низом или C препроцессору

-arch=sm_20 -- может пригодиться если использовать printf() в ядре так как эта функция появилась только в Fermi

-O -- имя бинарника

Пример простой программы на CUDA

/** Build in Linux:* export PATH=$PATH:/usr/local/cuda/bin* rm -f test_cuda_app* nvcc main.cu -o test_cuda_app* ./test_cuda_app)*//** Build in Windows:* :: Import Visual C++ 2013 toolchain, etc.* ::-----------------------------------------------------------------* set VS_DIR=%VS120COMNTOOLS:~,-15%\VC* set PATH=%PATH%;%VS_DIR%* call "%VS_DIR%/vcvarsall.bat" x86* ::-----------------------------------------------------------------* nvcc main.cu -o main

* nvcc main.cu -o main -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin"

*/

#include <algorithm>#include <malloc.h>#include <stdio.h>  __global__ void sum(float *z, float *x, float *y){   int i = blockIdx.x * blockDim.x + threadIdx.x;   z[i] = x[i] + y[i];}int main(void){   const int N = 1024;    float *x=0, *y=0, *z=0, *d_x=0, *d_y=0, *d_z=0;   x = (float*)malloc(N*sizeof(float));   y = (float*)malloc(N*sizeof(float));   z = (float*)malloc(N*sizeof(float));    cudaMalloc(&d_x, N*sizeof(float));   cudaMalloc(&d_y, N*sizeof(float));   cudaMalloc(&d_z, N*sizeof(float));    for (int i = 0; i < N; i++) {     x[i] = i * 1.0f;     y[i] = i * 2.0f;   }    // Args for cud memcpy: dst, src, direction   cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);   cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);    const int Blocks = 128;   sum<<<Blocks, N/Blocks>>>(d_z, d_x, d_y);   cudaMemcpy(z, d_z, N*sizeof(float), cudaMemcpyDeviceToHost);    float maxError = 0.0f;   for (int i = 0; i < N; i++)     maxError = max(maxError, abs(z[i] - (x[i] + y[i])));    printf("Max error: %f\n", maxError);}

Code snippets

Размер блока в элементах и размер грида с окрулением на верх. Вообще если количество элементов не обязательно кратно количеству потоков скорее всего придётся вставлять дополнительную проверку.

Вычисление размера сетки

dim3 block(elementsInBlocks);

dim3 grid( (totalElements + block.x - 1)/block.x );

kernelCall<<<grid,block>>> ();    // invoke kernel asynchronously. use "cudaDeviceSynchronize" to force waiting in all kernels

cudaDeviceSynchronize();              // wait for completion of kernel. Example of implicit sync is call cudaMemcpy

cudaDeviceReset();                          // reset device before end of application

Макрос для чеканья кода возврата

#define CHECK(call) \

{\

 const cudaError_t error = call; \

 if (error != cudaSuccess) \

 { \

 printf("Error: %s:%d, ", __FILE__, __LINE__); \

 printf("CUDA CODE:%d, reason: %s\n", error, cudaGetErrorString(error)); \

 exit(1); \

 } \

}

Установка текущего GPU устройства для вычислений

// set up device

int deviceCount = 0;

CHECK(cudaGetDeviceCount(&deviceCount));

int dev = 0;

cudaDeviceProp deviceProp;

CHECK(cudaGetDeviceProperties(&deviceProp, dev));

printf("Using Device %d: %s\n", dev, deviceProp.name);

CHECK(cudaSetDevice(dev));

Code ограничения

1. __global__ функции должны возвращать void

2. Имеют доступ только к глобальной памяти

3. Не поддерживается переменное количество параметров

4. Не поддерживаются статические переменные внутри функции

5. Не поддерживаются указатели на функции

6. На устройствах Fermi максимальное количество потоков на ThreadBlock - 1024. 

    А максимальный размер сетки [65535, 65535,65535]

Отладка

1. printf в Fermi (CUDA 6.0+)

2. запустить на <<<1,1>>> и посмотреть на поведение

3. nvprof ./binaryApplication - выдаст количество вызовов на функцию, простую статистику по времени работы. Время работы измеряется точнее чем через gettimeofday.  nvprof 

это comand line profiler (CUDA 5.0+)

nvprof --print-gpu-trace ./binaryApplication  - выдаст конфигурацию ядра, размер используемых регистров на поток и shared memory.

4. Take some metrics

Link to metrics: http://docs.nvidia.com/cuda/profiler-users-guide/index.html#metrics-reference

Required Global Memory Load Throughput - includes replayed memory load instructions that required more than one memory transaction

Requested Global Memory Load Throughput - does not include replays

gld_efficiency = Requested Global Memory Load/Required Global Memory Load

nvprof --metrics gld_throughput,gld_efficiency,gld_transactions ./binaryApplication  

4. nvidia-smi -L -- печатает GPU установленные в системе

(Path on Windows 10 to this application "C:\Program Files\NVIDIA Corporation\NVSMI\nvidia-smi.exe" )

5. nvidia-smi -q -i 0 -- печатает детали о GPU-0 установленной в системе

6. установить переменную окружения CUDA_VISIBLE_DEVICES

  CUDA_VISIBLE_DEVICES=2 - в этом случае драйвер NVIDIA сделает так что устройства 2 для вашего бинарника будет устройством 0.

  CUDA_VISIBLE_DEVICES=2,3 - в этом случае драйвер NVIDIA сделает так что видны только устройства 2,3 для рантайма.

7. nvvp - standalone visual profiler. Есть как отдельное приложение, так и является частью NSight.

8. Большинство счётчиков аппаратно находятся на SM, а не на весь GPU. Некоторые счётчики являются взаимоисключающими. Нет никаких гарантий, что при прочих равных 

счётчики будут показывать одно и тоже например из-за логики Warp Scheduler-а.

9. nvprof --metrics branch_efficiency ./testApp -- посмотреть процент веток который выполняется без divergence

10. nvcc -g -G -- форсировать компилятор не выполнять оптимизацию

11. -Xptxas -v,-abi=no -- увидеть статистику по потребляемым ресурсам на поток

12. -Xptxas -dlcm=ca -- отключить использование L1 кеша на доступ в глобальную память из ядер

13. -Xptxas -dlcm=cg -- включить использование L1 кеша на доступ в глобальную память из ядер

14. nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory" -- посмотреть частоту работы памяти

15. Build for several virtual architectures via embeding PTX code:

nvcc -gencode arch=compute_60,\"code=compute_60\" -gencode arch=compute_50,"code=compute_50"  test.cu 

15. Build for several virtual architectures via embeding PTX code and specific version of SASS:

nvcc 

-gencode arch=compute_60,\"code=compute_60\" 

-gencode arch=compute_50,"code=compute_50"  -gencode arch=compute_50,"code=sm_50"

Futher generated code can be inspected by cuobjdump.

Use nvprof to generate profile which will be observable in NVVP:

https://kth.instructure.com/courses/20917/pages/tutorial-nvvp-visualize-nvprof-traces

16. watch nvidia-smi

17. nvprof --events branch,divergent_branch ./simpleDivergence

(event counters for branch and divergent branch)

18.cudaGetLastError()

if multiple calls to cudaGetLastError return an error code, the calling application knows that each of those errors is distinct and different from the others (though their cause can still be related)

19.cudaDeviceSynchronize()

Blocks until the device has completed all preceding requested tasks. cudaDeviceSynchronize() returns an error if one of the preceding tasks has failed. If the cudaDeviceScheduleBlockingSync flag was set for this device, the host thread will block until the device has finished its work.

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DEVICE.html#group__CUDART__DEVICE_1g10e20b05a95f638a4071a655503df25d

20. Check kernel call

kernel_function<<<grid, block>>>(argument list);

CHECK(cudaDeviceSynchronize());

CHECK(cudaDeviceSynchronize()) blocks the host thread until the device has completed all preceding requested tasks, and ensures that no errors occurred as part of the last kernel launch. This technique should be used just for debugging purposes, because adding this check point after kernel launches will block the host thread and make that point a global barrier.

Опции профилировщика

1. nvprof --metrics archived_occupancy ./myapp args

The achieved occupancy of a kernel is defined as the ratio of the average active warps per cycle to the maximum number of warps

supported on an SM. A thread block is called an active block when compute resources, such as registers and shared memory, 

have been allocated to it. The warps it contains are called active warps. 

Active warps can be further classified into the following three types:

2. nvprof --metrics gld_throughput ./myapp args

The memory read efficiency of the kernel GB/s

3. nvprof --metrics inst_per_warp

The average number of instructions executed by each warp

3. nvprof --metrics gld_efficiency./myapp args

The ratio of requested global load throughput to required global load throughput.

Required Global Memory Load Throughput includes replayed memory load instructions that required more than one memory transaction, whereas Requested Global Memory Load Throughput does not include replays.

4. nvprof --metrics gld_transactions ./myapp args

The number of global load transactions.

5. nvprof --metrics branch_efficiency ./myapp args

The ratio of non-divergent branches to total branches

6. nvprof --metrics gld_efficiency,gst_efficiency ./myapp args

Memory load/store effi ciency metrics

Разные вещи

1. GPU L1 задизайнин только для пространственной локальности, не временной. 

    CPU L1 сделан и для временной и пространственной.

    Так что частое обращение по одной памяти для GPU никак не влияет на шанс остаться в L1 кеше [3, p.163]

( Uncached loads do not pass through the L1 cache and are performed at the granularity of memory segments (32-bytes) and not cache lines (128-bytes) [3,p.163])

2. Когда чтение с кешами  - кеши работают с гранулярностью чанков по 128 байт.

    Чтение без кешей работает с гранулярностью чанки по 32 байта. 

    Поэтому при невыровненном доступе иногда имеет смысл отключить кеширование.

3. Следует быть аккуратным в интерпретации метрики gld_efficiency - она работает по разному в зависимости от влюченного/выключенного кеша L1.

4. Не следует забывать, что запуск ядра асинхронен к коду CPU thread-а.

5. Уникальный id потока в пределах 3d блока - threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x [3, p.81]

6. CUDA_DEVICE_MAX_CONNECTIONS environment variable - adjust the number of concurrent hardware connections (requests from various cuda streams fall into this queue), up to 32, for a Kepler device. Default the number of concurrent hardware connections is limited to eight [3, p.284]

---------------------------------------------------------------------------------------------------------------------------------------------------

7. Опции комилятора

http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html

CUDA_DEVICE_MAX_CONNECTIONS environment variable 

Adjust the number of concurrent hardware connections, up to 32, for a Kepler device. 

Default the number of concurrent hardware connections is limited to eight [3, p.284]

9. Dynamic parallelism. Доступен для устройств с СС >= 3.5. Из потока GPU, выполняемого в рамках какого-то ThreadBlock можно запустить новую сетку. Эта сетка (grid) в принципе видна всем потокам в пределах этого блока. При компиляции кода с этой фишкой нужно компилировать код со следующими флажками:

Запускаемая сетка будет выполняться на том же устройстве. Максимально возможное количество рекурсивных вызовов 24.

10. Синхронизация:

  explicit: cudaDeviceSyncronize, cudaStreamSyncronize, cudaEventSyncronize

  implicit: cudaMemcpy and other memory relative functions [3, p.277]

11. Generate intermediate PTX file nvcc --ptx -o my.ptx source.cu

12. There are three major factors that can limit performance for a kernel:

 - Memory bandwidth

 - Compute resources

 - Instruction and memory latency

13. All advices to improve performance:

  1. Keep more concurrent warps active within an SM

  2. Assign more independent work to each thread/warp

  3. Optimize memory access patterns (maximize the use of bytes that travel on the bus)

  4. Sufficient concurrent memory accesses (hide memory latency)

14. #include <cuda_runtime_api.h> 

// include file for c/c++ source files to enable calling of functions like cudaMalloc

To call kernel functions from host C/C++ code compiled with usual compile you should create wrappers with "extern "C"" name mangling specification inside CUDA source file (.cu) and call from wrapper with <<<...>>> constructions kernel code.

15. cudaGetLastError - return last error and cleanup error state to cudaSuccess [3,p.437]

      cudaPeekLastError - return last error only

16. nvcc -G -g test.cu флажки для создания отладочной сборки

17. For use cuda-memcheck and racecheck: 

      nvcc -G -g test.cu.

cuda-memcheck -- detects array out-of-bounds errors, and misaligned device memory accesses. Very useful because such errors can be tough to track down otherwise

cuda-memcheck --tool racecheck -- this checks for shared memory race conditions

cuda-memcheck --tool initcheck -- detects reading of uninitialised device memory

But if it will lead to very slow binary in which you can not reproduce problem then use:

      nvcc -lineinfo -Xcompiler -rdynamic test.cu. for GCC and nvcc -lineinfo -Xcompiler /Zi test.cu for Visual Studio

18. Locking the clock

    nvidia-smi -i 0 -pm 1

    nvidia-smi -i 0 -ac MEM_CLOCK, COMPUTE_CLOCK

Чтение из глобальной памяти

[3,p.160]. Возможно будет идти через L1/L2 кеши.

Compiler flags: 

-Xptxas -dlcm=cg -- отключает использование L1 кеша. В этом случае варпы читают память чанками по 32 байта. В случае включенного L1 кеширования чанки памяти составляют 128 байт.

"On Kepler K10, K20, and K20x GPUs, the L1 cache is not used to cache global memory loads. The L1 cache is exclusively used to cache register spills to local memory"

Synchronization and memory issues

1. The order in which a GPU thread writes data to different memories, such as shared memory, global memory, page-locked host memory, or the memory of a peer device, is not necessarily the same order of those accesses in the source code.

2. void __syncthreads();

__syncthreads  acts as a barrier point at which threads in a block must wait until all threads have reached that point.

__syncthreads also ensures that all global and shared memory accesses made by these threads prior to the barrier point are visible to all threads in the same block

__syncthreads is used to coordinate communication between the threads of the same block

3. void __threadfence_block()

Ensures that all writes to shared memory and global memory made by a calling thread before the fence are visible to other threads in the same block after the fence. Memory fences do not perform any thread synchronization and so it is not necessary for all threads in a block to actually execute this instruction.

4. void __threadfence();

__threadfence stalls the calling thread until all of its writes to global memory are visible to all threads in the same grid.

5. void __threadfence_system()

Memory fence across the system (including host and device). It stalls the calling thread to ensure all its writes to global memory, 

page-locked host memory, and the memory of other devices are visible to all threads in all devices and host threads.

6. volatile qualifier

Declaring a variable in global or shared memory using the volatile qualifier prevents compiler optimization which might temporally cache data in registers or local memory. 

Therefore, any reference to this variable is compiled to:

-- Global memory read skiped cache

-- Global memory write instruction that skips the cache.

Features of different Compute Capability video-cards: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions

"CTA term" is the same as "thread block" term: https://docs.nvidia.com/cuda/parallel-thread-execution/#cooperative-thread-arrays

So CTA is hardware implementation of thread-block. What is really executed not grid, but really CTAs.

Approximate and rough chips naming schema:

**100 -- nominal

**102 -- each +2 in code name divide by two - available resources and physical size of chip

GPU/CUDA Libraries for Dense Lin.Algebra at BLAS level

1.  cuBLAS

2.  cuSOLVER

3.  MAGMA

4.  ArrayFire

GPU/CUDA Sparse lin.algebra at solvers level

1.  Iterative solvers: can be implemented with maxtrix-vector multiplication like cuSPARSE

2.  Direct solvers: cuSOLVER, superLU, RAL project (https://epubs.stfc.ac.uk/work/12189719)

2.  cuSOLVER

3.  MAGMA

4.  ArrayFire

Various display information

- Display device memory information only

nvidia-smi -q -i 0 -d MEMORY

- Display device utilization information only

nvidia-smi -q -i 0 -d UTILIZATION

Some extra terminology

pseudo-fp16 - store as fp16m but compute in fp32

Measure perfomance

Peak Bandwidth -- is theoretical memory bandwith. Can be estimated from characterstic of devices installed in system.

Estimated Transfer bytes / (stop-start) -- is practical bandwidth 

Peak TFLOPS -- is theoretical computaion bandwidth. Can be estimated from characterstic of devices installed in system.

Estimated TLOPS for execution / (stop-start) -- is practical bandwidth 

Some metric relative to performance

"flops_sp" gives a count of floating point operations 

nvprof --metrics flops_sp <path_to_binary>"

"gld_throughput" gives read efficiency of the kernel in GB/second */

https://docs.nvidia.com/cuda/profiler-users-guide/index.html#metrics-reference

http://on-demand.gputechconf.com/gtc/2013/presentations/S3046-Performance-Optimization-Strategies-for-GPU-Accelerated-Apps.pdf