应用的性能可能有 75% 都花费在内存相关问题上。
NVPROF and NVVP
这俩是调试工具，不知道是不是基于CUPTI (CUDA Profiler Tools Interface)。
NVPROF是命令行工具，nvvp是可视化工具。
nvvp有四个模块：Timeline，Summary，Guide，Analysis results
其中 Guide 适合新手，新手应该多注意。
为了分析出应用的性能瓶颈，我们需要timeling analysis 和 metric analysis
编译：
要include cuda_profiler_api.h 和 helper_functions.h
命令1
nvcc -I"D:/cuda/NVIDIA Corporation_v10.1/common/inc" -gencode arch=compute_61,code=sm_61 -Xcompiler /wd4819 -o sgemm sgemm.cu
命令2 创建 nvvp文件
nvprof -o sgemm.nvvp .\sgemm.exe
13388 NVPROF is profiling process 13388, command: .\sgemm.exe
Operation Time= 0.0046 msec
13388 Generated result file: D:\codes\Learn-CUDA-Programming-master\Learn-CUDA-Programming-master\Chapter02\02_memory_overview\01_sgemm\sgemm.nvvp
运行nvvp，打开 sgemm.nvvp文件
Global memory / device memory
global memory 是从host复制过去的默认空间
cudaMalloc 和 cudaFree cudaMemcpy
coalesced memory access / uncoalesced memory access
warp：32threads running in SMs。
例如 一个SM运行两个block，每个block128线程，则一共 8 个 warp。这些 warp 公用一个SM的shared memory。
每个warp中的线程以SIMT模式运行，也就是所有线程同时运行同一个命令
warp下的内存访问需要coalesced memory access，squential memory access is adjacent.
总而言之尽量连续访问、数据aligned
texture memory
一般用来存储大部分线程都要访问的少量数据，比如参数之类。同一个warp的线程访问同一个地址会导致所有线程在每个时钟周期请求数据，但是在texture memory上有优化。2D访问、3D访问也有优化，总之就是读的快、不能写。
shared memory
shared memory以banks的形式排列，要防止同一个warp里的线程同时访问同一个bank，至于不同warp之间的不太清楚。
registers
local variables are stored in registers
尽量别弄太多本地变量，寄存器和thread有关
pinned memory
尽量不在host 和 device之间传输数据
用pinned memory
把许多小数据传输合并成一个大batch
用计算来掩盖传输
传输数据时默认先创建一个pinned memory，然后将一个个page复制过去，再从pinned memory里将数据复制到device（通过DMA）。通过cudaMallocHost可以之间分配一个pinned memory。
Unified memory
给GPU和CPU提供一个统一的地址
cudaMallocManaged分配空间，不是立刻分配，下一次host第一个访问该空间就分配host，否则device。(first touch)
当device需要访问某个page，但是page正在host映射的空间里，这时会让host取消映射，transfer这个page到device的物理空间，device再映射这个page。
warp per page 技术，每个warp负责64K，一个page：
__global__ void init(int n, float *x, float *y) {
int lane_id = threadIdx.x & 31;//本线程是该warp的第几个线程
size_t warp_id = (threadIdx.x + blockIdx.x * blockDim.x) >> 5;//blockDim.x是一个block的线程数
size_t warps_per_grid = (blockDim.x * gridDim.x) >> 5;//
size_t warp_total = ((sizeof(float)*n) + STRIDE_64K-1) / STRIDE_64K;//pages
// if(blockIdx.x==0 && threadIdx.x==0) {
// printf("\n TId[%d] ", threadIdx.x);
// printf(" WId[%u] ", warp_id);
// printf(" LId[%u] ", lane_id);
// printf(" WperG[%u] ", warps_per_grid);
// printf(" wTot[%u] ", warp_total);
// printf(" rep[%d] ", STRIDE_64K/sizeof(float)/32);
// }
for( ; warp_id < warp_total; warp_id += warps_per_grid) {
#pragma unroll//告诉编译器任意展开并行等都是安全的
for(int rep = 0; rep < STRIDE_64K/sizeof(float)/32; rep++) {
size_t ind = warp_id * STRIDE_64K/sizeof(float) + rep * 32 + lane_id;
if (ind < n) {
x[ind] = 1.0f;
// if(blockIdx.x==0 && threadIdx.x==0) {
// printf(" \nind[%d] ", ind);
// }
y[ind] = 2.0f;
}
}
}
}data prefetching: cudaMemPerfetchAsync()
预先告诉程序下一个设备是谁。这个最快。
cudaMemAdvice()系列API：
SetReadMostly、PreferredLocation、SetAccessBy等。
UM技术把device和host的内存看作同一个内存，可以用来解决很多GPU内存不够的问题。
