1
2
3
4
5
6
7
8
9
10
11
12
13
|
__global__ void gpu_matrix_transpose(int in[N][M], int out[M][N])
int x = threadIdx.x + blockDim.x * blockIdx.x;
int y = threadIdx.y + blockDim.y * blockIdx.y;
if( x < M && y < N)
out[x][y] = in[y][x];
1
2
3
4 | # nvcc（NVIDIA C编译器）用于把内核编译成PTX格式
# Parallel Thread Execution (PTX) 是一个低级虚拟机和指令集
# graphics driver 将PTX转换成可执行的二进制代码（SASS）
nvcc -o t1866 t1866.cu
register , local memory 每个线程独享的 （快）
shared memory/L1 block线程共享 （快，小）共享内存实际上是在 GPU 芯片上实现的，因此我们可以称之为片上内存。由于它是片上的，其速度也比全局内存快得多。通常最多可达48KB共享（或64KB、96KB等）。延迟非常低。吞吐量非常高：>1 TB/s的总和。
1
2
3
4
5
6
7
8
9
10
11 | __global__ void gpu_matrix(int* a, int* b, int* c, int m, int n, int k)
// share
__shared__ int sub_a[BLOCK_SIZE][BLOCK_SIZE];
__shared__ int sub_b[BLOCK_SIZE][BLOCK_SIZE];
。。。。 写操作
// 线程并行执行，要先写后读 同一块（block）中的线程之间同步执行
__syncthreads();
。。。。 读操作
线程间的交流通道
可编程的 cache
通过缓存数据减少 glabal memory 访存次数。
cuda bank conflict memory padding
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21 | __global__ void gpu_shared_matrix_transpose(int in[N][M], int out[M][N])
int y = threadIdx.y + blockDim.y * blockIdx.y;
int x = threadIdx.x + blockDim.x * blockIdx.x;
__shared__ int ken[BLOCK_SIZE+1][BLOCK_SIZE+1];//ken[32] warp 通过padding解决 bank conflict问题
if(x < M && y < N)
ken[threadIdx.y][threadIdx.x] = in[y][x];
__syncthreads();
int x1 = threadIdx.x + blockDim.y * blockIdx.y;
int y1 = threadIdx.y + blockDim.x * blockIdx.x;
if(x1 < N && y1 < M)
out[y1][x1] = ken[threadIdx.x][threadIdx.y];//32 bank
当您每个线程存储（或加载）超过 4 个字节时，即每个 warp 超过 128 个字节时，GPU 不会发出单个事务。最大事务大小为 128 个字节。每个事务的宽度为 128 个字节。bank conflicts是按事务进行的，而不是按请求、warp或指令进行的。
CPU只能访问到以下内存
global memory 所有线程共享 (慢，大) 物理实现，通常实现在 GPU 的动态随机荐取存储器(DRAM)中。这并非位于 GPU 芯片本身，而是由多个高速内存设备组成。这些设备与 GPU 相连。 延迟高（数百个周期）。吞吐量：高达约900 GB/s（Volta V100）。所有线程以及主机（CPU）都可以访问。
常量和纹理 memory （只读，相对global快）
CUDA 编程模型还假设主机和设备都在 DRAM 中维护各自的内存空间，分别称为Host内存和 Decice内存。因此，程序通过调用 CUDA 运行时来管理内核可见的全局、常量和纹理内存空间。这包括设备内存分配和释放以及主机和设备内存之间的数据传输。
Global 内存
默认模式。
尝试在L1中命中，然后是L2，再然后是全局内存（GMEM）。
加载粒度为128字节行。
使用 -Xptxas -dlcm=cg 选项编译 nvcc。
尝试在L2中命中，然后是全局内存（GMEM）。
不要在L1中命中，如果已经在L1中，则使该行失效。
加载粒度为32字节。
使L1失效，写回L2。
内存操作是按warp（32个线程）发出的。 段是32字节
与所有其他指令一样。
warp中的线程提供内存地址。
确定所需的行/段。
请求所需的行/段。
Warp请求32个不对齐的连续4字节的单元
int c = a[idx]; or int c = a[rand()%warpSize];(随机落到一个范围里)
没有任何额外请求未被swarp中的某个线程所使用, 总线利用率100%，每个请求的享节均得到利用。（pefect Coalescing 完美合并）
32个线程，提供32个地址，而内存控制器需要决定如何处理这种请求，它会讲这些地址根据其所在的行或者段进行合并或者分组（合并）。随后它不再请求32个地址单独地址，而是一组连续的行或者段。
如果实现了完全合并，则意味着我所描述的情况得以实现。所有地址均能通过最少量的线或段得到满足。
现在，这些数字，目即这些索引，都是相邻的。但它们相对于前一个示例有所偏移。实际上，它们会跨越一条线或线段的边界。
在这种情况下，它们跨越了边界，因此肉存控制器会将这些地址合并为两个独立的组，一组属于第一个边界区域，即内存地址从0 到 128 的部分，而太部分线程将归并到第二个区域，即内存地址从128到 256 的范围。因此，内存控制器需要请求两条行或多个段来处理此次请求。
如果我们请求两条缓存行，意味着内存控制器将需要检索 256字节。请记住，你不能要求少示最小量即一条缓存行或一段内存。若受限于从内存检索数据的能力，即内存瓶颈，性能将直接下降50%
全局内存优化指南
努力实现完美的合并（Coalescing）：
（对齐起始地址 - 可能需要填充）。
一个warp应在连续区域内进行访问。
确保有足够的并发访问以饱和总线：
每个线程处理多个元素。
多个加载可以被流水线处理。
索引计算通常可以被重用。
如果是按bank遍历，共享内存能够以最短的时间完成该加载操作。
如果在同一加载事务中请求，假设线程0需要字节0、1、2、3 ，而线程1需要字节 128、129、130、131。这是列式访问，若采用列式访问模式，共享内存将使这些事务串行化。共享内存会表示，在第一个周期，它将为第一个线程提供服务，在第二个周期，我将处理第二个线程，以此类推，直到所有在加载操作中向零号存储体请求项目的线程，即该指令在全 warp范围内发出的所有线程。
如果32个线程都访问同一个bank，共享性能最差，
bank conflicts
即两个或多个线程请求同一bank(或可说在共享内存的同一列)中的项。
统一内存 Unified Memory
统一内存提供托管内存来连接主机和设备内存空间。托管内存可作为具有公共地址空间的单一、连贯的内存映像从系统中的所有 CPU 和 GPU 访问。旨在简化编程模型，特别是简化内存模型。
我们并未改变这样一个事实， data DtoH and HtoD 这两个步骤是必须的。使用这个，只是简化编程工作。
cuda runtime是一个执行实体，它基于代码执行 DtoH and HtoD 。所以cuda runtime会关注编写的代码，并影响数据的移动（数据迁移，将数据迁移到需要它的处理器上面，以高效的runtime处理）。这种数据迁移机制：确保数据一次仅有一个处理器访问，保证全局一致性
1
2
3
4
5
6
7
8
9
10 | #define M 1000
#define N 500
#define K 1000
// 统一内存 直接从设备或主机功能中读取或写入其地址。 替代cudaMalloc,cudaMemcpy,cudaFree等 不需要分别申请Host内存和 Decice内存
__managed__ int a[M*N];
__managed__ int b[N*K];
__managed__ int c_gpu[M*K];
__managed__ int c_cpu[M*K];
// cudaMemcpy 已经包含类似 cudaDeviceSynchronize 的功能 阻塞等待所有cuda活动完成，在复制动作开始前
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23 | #include <stdio.h>
__managed__ float managedData[256]; // 声明统一内存
__global__ void myKernel() {
int tid = threadIdx.x;
managedData[tid] = tid * 2.0f; // 在 GPU 中修改
int main() {
// 启动内核
myKernel<<<1, 256>>>();
// 等待 GPU 完成
cudaDeviceSynchronize();
// 在 CPU 中访问
for (int i = 0; i < 256; i++) {
printf("managedData[%d] = %f\n", i, managedData[i]);
return 0;
数据迁移原理
数据迁移通过page fault 触发：
page fault on the GPU: data H to D
page fault on CPU： data D to H
当该页未存在于 GPU上或未驻留时，便发生了页面错误，导致数据迁移。并在GPU内存中物理实例化该页。此后，若发生页面错误，任何触及该页的额外代码将不再引发减速，没有任何开销。它只是以正常的速度继续进行。因此，页面错误通常仅在首次有CUDA 线程访间该页面时发生一次，随后访问该页面的线程通常不会观察到页面错误。
若我以高频率访问代码，无论是在 CPU或者GPU上。此页面将会来回切换。
__managed__ 和 cudaMallocManaged 都用于在 CUDA 中实现统一内存，但它们的使用方式和上下文有所不同。
__managed__
用途：用于声明统一内存变量，使得这些变量在 CPU 和 GPU 之间共享。
声明方式：在变量声明时直接使用 __managed__ 关键字。
作用范围：适用于全局变量或静态变量。
1 | __managed__ float managedData[256];
cudaMallocManaged
用途：用于动态分配统一内存，允许在运行时分配内存。
调用方式：使用 cudaMallocManaged() 函数进行分配。
作用范围：适用于动态分配的内存，通常用于堆内存。
1
2 | float *managedPtr;
cudaMallocManaged(&managedPtr, size * sizeof(float));
声明方式：
__managed__ 在变量声明时使用。
cudaMallocManaged 在运行时动态分配内存。
使用场景：
__managed__ 适合于全局、静态变量。
cudaMallocManaged 更灵活，适合于需要动态分配的内存。
内存管理：
__managed__ 变量的生命周期与程序相同。
cudaMallocManaged 需要在使用完后调用 cudaFree() 来释放内存。
// be careful with size type:
size_t size = 64ULL*1024*1024*1024;
cudaMallocManaged(&data, size);
作为内核代码，假设触及此分配时，任何尚未物理驻留的页面将通过需求分页系统被引入到 GPU 内存中驻留。现在，当我通过该进程耗尽或超出 GPU 内存时，某些方面必须做出让步。问题在于，那些最近未被访问过的页面将会被驱逐。它们仍然存在于 GPU 内存映射中，只是物理上不驻留在设备上。因此，若代码再次访问被置换的页面，将导致该页面重新迁移回 GPU。
代价：延迟会增加
CPU 和 GPU 对同一分配的访问
1
2
3
4
5
6
7
8
9
10
11 | __global__ void mykernel(char *data) {
data[1] = ‘g’;
void foo() {
char *data;
cudaMallocManaged(&data, 2);
mykernel<<<...>>>(data);
// no synchronize here
data[0] = ‘c’;
cudaFree(data);
因为运算是异步的，无法判断那个会先发生。所以必须要程序员解决顺序问题
系统级原子操作
1
2
3
4
5
6
7
8
9
10
11
12 | __global__ void mykernel(int *addr) {
atomicAdd_system(addr, 10);
void foo() {
int *addr;
cudaMallocManaged(addr, 4);
*addr = 0;
mykernel<<<...>>>(addr);
// cpu atomic:
__sync_fetch_and_add(addr, 10);
适用于多处理器
系统级原子操作使您能够以原子方式从所有处理器访问托管分配
统一内存自动做
类似于深拷贝的情况
编写复制操作的代码很复杂
统一内存使得这一过程变得简单
1
2
3
4
5
6
7
8
9
10
11
12
13 | void launch(dataElem *elem, int N) { // an array of dataElem
dataElem *d_elem;
// Allocate storage for array of struct and copy array to device
cudaMalloc(&d_elem, N*sizeof(dataElem));
cudaMemcpy(d_elem, elem, N*sizeof(dataElem), cudaMemcpyHostToDevice);
for (int i = 0; i < N; i++){ // allocate/fixup each buffer separately
char *d_name;
cudaMalloc(&d_name, elem[i].len);
cudaMemcpy(d_name, elem[i].name, elem[i].len, cudaMemcpyHostToDevice);
cudaMemcpy(&(d_elem[i].name), &d_name, sizeof(char *), cudaMemcpyHostToDevice);
// Finally we can launch our kernel
Kernel<<< ... >>>(d_elem);}
采用逐页处理的方式，即零散进行，其开销将远高于批量移动数据时的成本。
所以对于大数据量移动或者访问大块显存，使用cuda memcopy效率更好
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15 | __global__ void kernel(float *data){
int idx = …;
data[idx] = val;
int n = 256*256;
int ds = n*sizeof(float);
float *data;
cudaMallocManaged(&data, ds);
// to GPU
cudaMemPrefetchAsync(data, ds, 0);
Kernel<<<256,256>>>(data);
// to CPU
cudaMemPrefetchAsync(data, ds,cudaCpuDeviceId); // copy back to host
可以将数据预取到任何 GPU 或 CPU，这种灵活性有助于优化多设备之间的数据管理。
cudaMemPrefetchAsync 是异步的，不会阻塞主机线程。这意味着可以同时进行其他计算或内存传输。
提前预取数据可以减少后续访问时的延迟，尤其是在数据访问模式已知的情况下，能显著提高性能。专注于数据预取，为后续使用准备数据，而不是立即进行复制。
Advise runtime on expected memory access behaviors with:
cudaMemAdvise(ptr, count, hint, device);
Hints:
cudaMemAdviseSetReadMostly: Specify read duplication , 协商（不强求）只读，保证多处理数据一致，会复制多个副本
一旦某个处理器进行写入，即违反了您的提示，统一内存子系统将介入并使所有相关数据失效，迫使再次进行迁移，以确保数据一致性。
如果准守约定：每个处理器都拥有其本地副本，且不存在任何迁移。若处理器首次触及只读数据，此时将发生迁移，但此后该数据将永久保留，即便其他进程也在读取它。除了首个接触的处理器，其他处理器都可以全速访问。
cudaMemAdviseSetPreferredLocation: suggest best location
当首个 GPU 触及该数据时。数据迁移即刻发生，便在该 GPU 上驻留。第二块 GPU 接触到该数据时，不会将数据迁移至第二块，UM 系统将尝试建立映射，通过处理器间总线（NVLink/PCIe）处理读写请求。
cudaMemAdviseSetAccessedBy: suggest mapping
指定GPU通过映射方式访问数据，而不是数据迁移。映射比普通访问慢，但映射使得你能够通过总线访问它，无需通过页面错误来处理访问请求。若你知晓某个处理器偶尔会访问你确认位于另一处理器上的数据，那么提供这种通过集合访问的额外提示或许是有益的。
CUDA 编程模型通过异步编程模型为内存操作提供加速。异步编程模型定义了异步操作相对于 CUDA 线程的行为。
线程协同工作：一种通过通信（通过shared memory，warp shuffle），另一种方式则是通过同步。
规约：多个输入，一个输出
Atomic Functions
*c += a[i];
在GPU编程里面做不到，编译器将这些代码转换成实际可执行的指令
1
2
3
4
5
6 | Actual code the GPU executes:
LD R2, a[i] (thread independent) # 特定线程从输入数组中加载其值,我们称之为a[i ]，并将该值存入寄存器
LD R1, c (READ)
ADD R3, R1, R2 (MODIFY)
ST c, R3 (WRITE) # 将R3的值存入指针c的位置
当我们跨多个线程执行此操作时。观察到的现象是所有线程都正确完成了第1 步指令。但是后面的步骤：线程们实际上是在相互踩踏，
因为每个线程都在尝试执行或者同时执行，CUDA不会自动为您完成排序。如果代码的正确性依赖于线程执行的顺序，则该代码本身并不正确。
这个时候需要原子操作，将3步转成1步。实际上，它表现为我们所谓的“归约“操作，但归约是机器代码级别上的原子指令形式。在此实现中，不允许其他线程或任何其他活动访问由指针C所指示的位置。让线程间的行为串行化，但这样能确保行为定义明确、可预测且符合预期
原子硬件实际上是在一个名为L2 缓存的地方实现的。当线程发起原示指令时，L2缓存中存在一个协调执行机制，作为协调者它会逐一处理所有这些原子操作。
当多个线程尝试使用原子换作更新同一位置，线程操作的串行化可能导致性能影响。所以，我们通常不应期望原子操作，能以与普通内存操作相同的速率进行
原子操作所适用的数据类型及其在不同 GPU 架构上的影响范围可能会有所不同。
int my_position = atomicAdd(order, 1);
可用于确定下一个工作项、队列槽等。确定我在顺序中的位置
大多数原子操作返回一个值，该值是接收原子更新的位置的“旧”值。
预留缓冲区中的空间
希望树的顶层所有操作，必须在任何低级树操作开始之前完成。
由于这种同步作用于整个树的宽度，类似于作用于整个网格的宽度，我们将其称为全局同步。跨越block级别，作用于grid的线程同步。
把任务用kernel切分，kernel会按序完成，当kernel完成就意味着所有线程都完成了，kernel的启动边界提供一个全局同步屏障。
kernel启动本身存在一定开销，虽然相对于处理时间而言。这一开销通常较低，但如果内核执行大量工作，启动开销可能就不容忽视
与内核启动相关联的网格变小了，因为随着我沿树向下移动。树的宽度在缩小。
block排空法：kernel结束后，block是一个个退出的，完成一些收尾的工作。通过使用原子操作，跟踪这个活动，知道哪个block是最后完成的。我们就可以将额外的工作给它，因为我们知道其他所有线程块均已完成。cuda sameple threadFenceReduction
协作组：CUDA编程模型允许使用相当粗粒度的结构进行分解。主要是解决：CUDA在提供线程协作和线程分解、线程组分解方面的构造或原语不够丰富的问题。协作组提供了一套新的内置函数和基本组件，使我们能够构建规模更为灵活的线程组，这些线程组能够协同工作，共同执行任务。
1
2
3
4
5
6 | for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
__syncthreads(); // outside the if-statement
希望创建的核函数能够实现一种解耦，即核函数的规模（换言之，网格的大小，即执行操作的线程数量）与数据集大小之间的解耦
能够有效加载和操作任意数据大小的kernel，将初始时对输入数据集大小进行处理，该大小与网格的宽度相对应。
固定数量的线程，等于网格的宽度。即预先确定线程数量。
1
2
3
4
5
6
7 | // block id
int idx = threadIdx.x+blockDim.x*blockIdx.x; // block width
while (idx < N) {
sdata[tid] += gdata[idx];
// global id
idx += gridDim.x*blockDim.x; // grid width
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16 | __global__ void reduce_a(float *gdata, float *out){
__shared__ float sdata[BLOCK_SIZE];
int tid = threadIdx.x;
sdata[tid] = 0.0f;
size_t idx = threadIdx.x+blockDim.x*blockIdx.x;
while (idx < N) { // grid stride loop to load data
sdata[tid] += gdata[idx];
idx += gridDim.x*blockDim.x;
for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
__syncthreads();
if (tid < s) // parallel sweep reduction
sdata[tid] += sdata[tid + s];
if (tid == 0) atomicAdd(out, sdata[0]);
warp shuffle
允许Warp内部实现这种直接的线程间通信。
warp 由32个线程组成，同步运行
1
2
3
4 | __shfl_sync(): 从任意 lane ID 复制（任意模式）lane 是warp 里面的thread索引
__shfl_xor_sync():
__shfl_up_sync():
__shfl_down_sync(unsigned mask, float value, unsigned delta, unsigned width=warpSize): 将同一个 warp 中某个线程的值传递给其下delta/offset 个线程
使用同步“mask”掩码用于指定哪些线程参与操作
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33 | __global__ void reduce_ws( float *gdata, float *out )
__shared__ float sdata[32];
int tid = threadIdx.x;
int idx = threadIdx.x + blockDim.x * blockIdx.x;
float val = 0.0f;
unsigned mask = 0xFFFFFFFFU; // 16进制32 位中的每一位都是 1无符号整数 所有线程都参与
int lane = threadIdx.x % warpSize; // 一个线程在其所在的 warp 中的索引
int warpID = threadIdx.x / warpSize; // 当前线程所在的 warp 的索引
while ( idx < N ) /* grid stride loop to load */
val += gdata[idx]; // 可以不用share memory
idx += gridDim.x * blockDim.x;
/* 1st warp-shuffle reduction */
for ( int offset = warpSize / 2; offset > 0; offset >>= 1 )
val += __shfl_down_sync( mask, val, offset );
if ( lane == 0 )
sdata[warpID] = val;
__syncthreads(); /* put warp results in shared mem */
/* hereafter, just warp 0 */
if ( warpID == 0 )
/* reload val from shared mem if warp existed */
val = (tid < blockDim.x / warpSize) ? sdata[lane] : 0;
/* final warp-shuffle reduction */
for ( int offset = warpSize / 2; offset > 0; offset >>= 1 )
val += __shfl_down_sync( mask, val, offset );
if ( tid == 0 )
atomicAdd( out, val );
减少了每个线程块所需的共享内存量
wrap shuffle 操作机器码上面是单一指令，减少了指令数量
减少了显式同步的频率，相较于共享内存扫描，后者在每次循环迭代中每个扫描操作都需要进行一次同步线程
将单一值广播至整个线程束中的所有线程，仅需一条指令
Cooperative Groups 协作组
即在多个执行单元间实现基本合作，Cooperative Groups可以做到全Grid同步
__syncthreads()充当block级别同步屏障，而Cooperative Groups有一个抽象概念thread group：使得一组线程可以通信和同步
thread block： 启动的线程块中所有线程的集合。
32,4 这些值必须小于或等于32，且2必须的幂
1
2
3
4
5
6 | // 获取当前线程块的线程组。
thread_block g = this_thread_block();
// 将上面的thread block 分成每份有32个线程的tiles 在协作组中,没有明确的线程束warp概念。但是可以理解成tiles对应warp
thread_group tile32 = tiled_partition(g, 32);
// 分解成大小为4的tiles
thread_group tile4 = tiled_partition(tile32, 4);
thread block
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30 | // Per-Block
g = this_thread_block();
reduce(g, ptr, myVal);
// Per-Warp
g = tiled_partition(this_thread_block(), 32);
reduce(g, ptr, myVal);
// __device__ 修饰符表明该函数或变量只能在 GPU 上被调用或访问，不能直接从主机代码（CPU）中访问。只能在核函数中调用
__device__ int reduce(thread_group g, int *x, int val) {
int lane = g.thread_rank(); // 返回当前线程在线程组中的索引。
for (int i = g.size()/2; i > 0; i /= 2) {
x[lane] = val; g.sync(); // 确保所有线程在进行下一步操作前都完成了当前步骤，避免数据竞争。
if (lane < i) val += x[lane + i]; g.sync();
return val;
g = tiled_partition<16>(this_thread_block());
tile_reduce(g, myVal);
template <unsigned size>
__device__ int tile_reduce(thread_block_tile<size> g, int val) {
for (int i = g.size()/2; i > 0; i /= 2) {
val += g.shfl_down(val, i);
return val;
1
2
3
4
5
6
7 | __global__ kernel() {
grid_group grid = this_grid();
// load data
// loop - compute, share data
grid.sync();
// device wide execution barrier
协同启动内核要求
grid size不能超过一定大小，block没有分配到SM, 会造成死锁。一旦block被分配至 SM，它将永久留至任务完成。
multi grid group
1
2
3
4
5
6
7 | __global__ void kernel() {
multi_grid_group multi_grid = this_multi_grid();
// load data
// loop - compute, share data
multi_grid.sync();
// devices are now synced, keep on computing
Coalesced group 合并的线程
一个线程束是一组同步执行的线程集合（SIMD）。当向 warp 中的一个线程发出指令时，该指令也会同时发给 warp 中的所有其他线程。
一旦理解了线程束（warp）的概念，我们可能会问，当存在条件代码导致单个线程束内的线程在条件行为上产生分歧时，存在一个执行引擎，负责处理这种所謂的分岔状态或分岔，即线程束分岔（warp divergence）。它使得一些线程能够遵循“如果”路径，一些线程遵循“那么”路径，还有一些线程遵循“否则”路径，而无需深入细节。存在一个引擎，允许分歧行为。
当我们想象一些线程遵循某一条执行路径，而另一条线程遵循另一条执行路径时，我们能快速意识到，存在某些情况，我们可能无法拥有一个完整的线程束，或者可能无法拥有一个完整的线程束。
Coalesced group就是告诉我那些线程正在同步执行。
1
2
3
4
5
6
7
8
9
10 | inline __device__ int atomicAggInc(int *p)
coalesced_group g = coalesced_threads();
int prev;
if (g.thread_rank() == 0) {
prev = atomicAdd(p, g.size());
prev = g.thread_rank() + g.shfl(prev, 0);
return prev;
PINNED (NON-PAGEABLE) MEMORY 页锁内存
主机端存在虚拟内存，主机内存不足是会将内存数据交换到虚拟内存中，虚拟内存就是主机中的磁盘空间，需要该页时再重新从磁盘加载回来。这样做可以使用比实际内存更大的内存空间。
函数cudaMalloc()将分配标准的，可分页的主机内存。
cudaHostAlloc()将分配页锁定的主机内存。页锁定的主机内存也称为固定内存或不可分页内存，
它的重要属性就是:操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全的使用应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位
设备内存与锁页内存之间的数据传输可以与内核执行并行处理，方便多流计算
锁页内存可以映射到设备内存，减少设备与主机的数据传输。
在前端总线的主机系统锁页内存与设备内存之间的数据交换会比较快;并目可以是write-combining的，此时带宽会很大。
cudaMalloc 与 cudaMallocHost 的区别总结
内存位置：
cudaMalloc：在设备（GPU）上分配内存，适用于 GPU 计算。
cudaMallocHost：在主机（CPU）上分配页锁定内存，主要用于 CPU 与 GPU 之间的数据传输。
访问方式：
cudaMalloc：分配的内存只能由 GPU 直接访问，CPU 需要通过 cudaMemcpy 进行数据传输。
cudaMallocHost：分配的内存可以被 GPU 快速访问，适合高效的数据传输。
cudaMallocHost：页锁定内存提高了数据传输性能，但不用于计算。
cudaMalloc：设备内存用于 GPU 执行计算。
释放方式：
cudaFree：释放通过 cudaMalloc 分配的内存。
cudaFreeHost：释放通过 cudaMallocHost 分配的内存。
cudaMalloc 主要用于 GPU 计算，而 cudaMallocHost 则用于主机内存，适合高效的数据传输。
两者各有特定用途，不能相互替代。使用 cudaMalloc 进行计算，使用 cudaMallocHost 进行快速数据传输。
页锁内存（Pinned Memory）的缺点
内存限制：
页锁内存的使用量通常受到系统限制（上限物理内存总量），过多使用可能导致系统内存不足，影响主机性能。
性能开销：
尽管页锁内存可以提高数据传输性能，但在某些情况下，使用它可能导致系统整体性能下降，尤其是在 CPU 进行大量内存操作时。
资源占用：
页锁内存会占用系统内存的页表项，可能影响其他应用程序的性能。
分配和释放开销：
分配和释放页锁内存的开销通常比普通内存更高，可能导致性能下降。
不适合频繁分配：
由于分配和释放成本较高，频繁使用页锁内存可能不利于性能。
每个stream中的操作按顺序执行，但不同stream中的操作可以并行执行，形成grid级别的并行。
使用多个stream可以重叠计算（kernel）和数据传输(cudaMemcpyAsync)，从而隐藏内存访问延迟。时间上重叠
CUDA流中排队的操作和主机都是异步的，所以排队的过程中并不耽误主机运行其他指令，所以这就隐藏了执行这些操作的开销。
我们主要的并发需求是同时调度从主机到设备以及从设备到主机的数据复制操作。
1
2
3
4
5
6
7
8
9
10
11
12 | cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream1); // potentially overlapped
// 第三个参数实际上是动态分配的共享内存大小，该内存将提供给内核。
// 第四个参数省略 称为空流 （null stream）
kernel<<<grid, block, 0, stream2>>>(...);
cudaStreamQuery(stream1); // test if stream is idle
cudaStreamSynchronize(stream2); // force CPU thread to wait
cudaStreamDestroy(stream2);
这是深度优先：即我们在完全横向展开操作的宽度之前，先沿着深度方向（三阶段按序入流）进行。
如果三个阶段，按阶段分stream，按块入流。广度优先
不指定流（或使用0作为流）的内核或 cudaMemcpy 正在使用默认流。当您未使用显式流，也未利用流 API 的任何特性时，所有工作均在默认流中执行。
首先，默认流将同步执行，这意味着它会强制所有先前发出的 CUDA 活动（无论在哪个流中发出）完成，然后默认流发出的项才会执行。（红色部分内）
此外，默认流要求在发出此默认流项之后发出的任何其他活动，必须等到默认流项完成后方可开始。
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87 | int main()
cudaDeviceProp prop;
int whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDeviceProperties(&prop, whichDevice);
// 是否支持多流
if( !prop.deviceOverlap )
printf("Your device will not support speed up from multi-streams\n");
return 0;
cudaEvent_t start, stop;
float elapsedTime;
// 定义三个流 不一定越多越好，看硬件
cudaStream_t my_stream[3];
int *h_a, *h_b, *h_c;
int *d_a0, *d_b0, *d_c0;
int *d_a1, *d_b1, *d_c1;
int *d_a2, *d_b2, *d_c2;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// 创建
cudaStreamCreate(&my_stream[0]);
cudaStreamCreate(&my_stream[1]);
cudaStreamCreate(&my_stream[2]);
cudaMalloc((void**) &d_a0, N * sizeof(int));
cudaMalloc((void**) &d_b0, N * sizeof(int));
cudaMalloc((void**) &d_c0, N * sizeof(int));
cudaMalloc((void**) &d_a1, N * sizeof(int));
cudaMalloc((void**) &d_b1, N * sizeof(int));
cudaMalloc((void**) &d_c1, N * sizeof(int));
cudaMalloc((void**) &d_a2, N * sizeof(int));
cudaMalloc((void**) &d_b2, N * sizeof(int));
cudaMalloc((void**) &d_c2, N * sizeof(int));
cudaHostAlloc((void**) &h_a, FULL_SIZE * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**) &h_b, FULL_SIZE * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**) &h_c, FULL_SIZE * sizeof(int), cudaHostAllocDefault);
for(int i = 0; i<FULL_SIZE; i++)
h_a[i] = rand()%1024;
h_b[i] = rand()%1024;
cudaEventRecord(start);
for(int i = 0; i < FULL_SIZE; i += N * 1)
cudaMemcpyAsync(d_a0, h_a+i, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[0]);
//cudaMemcpyAsync(d_a1, h_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[1]);
//cudaMemcpyAsync(d_a2, h_a+i+N+N, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[2]);
cudaMemcpyAsync(d_b0, h_a+i, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[0]);
//cudaMemcpyAsync(d_b1, h_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[1]);
//cudaMemcpyAsync(d_b2, h_a+i+N+N, N*sizeof(int), cudaMemcpyHostToDevice, my_stream[2]);
kernel<<<N/256, 256, 0, my_stream[0]>>>(d_a0, d_b0, d_c0);
//kernel<<<N/256, 256, 0, my_stream[1]>>>(d_a1, d_b1, d_c1);
//kernel<<<N/256, 256, 0, my_stream[2]>>>(d_a2, d_b2, d_c2);
cudaMemcpyAsync(h_c+i, d_c0, N*sizeof(int), cudaMemcpyDeviceToHost, my_stream[0]);
//cudaMemcpyAsync(h_c+i+N, d_c0, N*sizeof(int), cudaMemcpyDeviceToHost, my_stream[0]);
//cudaMemcpyAsync(h_c+i+N+N, d_c0, N*sizeof(int), cudaMemcpyDeviceToHost, my_stream[0]);
cudaStreamSynchronize(my_stream[0]);
cudaStreamSynchronize(my_stream[1]);
cudaStreamSynchronize(my_stream[2]);
// 0流 默认流
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime , start, stop);
printf("Time: %3.2f ms\n", elapsedTime);
// cudaFree
return 0;
multi GPU device manager
多个主机线程可以共享一个设备，单个主机线程可以管理多个设备
1
2
3
4
5
6
7 | cudaGetDeviceCount(int *count)
cudaSetDevice(int device)
cudaGetDevice(int *device)
cudaGetDeviceProperties(cudaDeviceProp *prop, int device)
cudaSetDevice(i) 选择当前设备
cudaMemcpyPeerAsync(...) 用于点对点复制,实现设备间复制
cuda smaple deviceQuery
流（Streams）和事件（cudaEvent）具有隐式/自动的设备关联
cudaStreamWaitEvent() 可以用来同步不同设备的流，而 cudaEventQuery() 可以检查一个事件是否“完成”。
如果你在一个跟当前设备没关联的流里启动内核，那就会失败。
每个设备都有自己独特的默认流。
1
2
3
4
5
6
7 | cudaSetDevice(0);
cudaStreamCreate(&stream0); // 关联到设备0
cudaSetDevice(1);
cudaStreamCreate(&stream1); // 关联到设备1
Kernel<<<b, t, 0, stream1>>>(...); // 这些内核有可能
cudaSetDevice(0);
Kernel<<<b, t, 0, stream0>>>(...); // 同时执行
设备之间数据复制
系统拓扑支持，数据可以直接从一个设备复制到另一个设备，使用的是像PCIE或NVLink这样的连接。避免经过主机内存。
两个设备放入一个对等关系（“clique”）中，两种传输方向启用“peering”功能，使用GPUDirect P2P 传输，但是在同一对等组中放置的设备数量存在限制（8~9）
1 | cudaDeviceEnablePeerAccess(peerDevice, flags);
peerDevice：指定要访问的设备 ID。
flags：通常设为 0，表示没有特殊要求。
1
2
3
4
5
6
7
8
9 | cudaSetDevice(0);
cudaDeviceCanAccessPeer(&canPeer, 0, 1); // test for 0, 1 peerable
cudaDeviceEnablePeerAccess(1, 0); // device 0 sees device 1 as a “peer”
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0, 0); // device 1 sees device 0 as a “peer”
cudaMemcpyPeerAsync(dst_ptr, 0, src_ptr, 1, size, stream0); //dev 1 to dev 0 copy
cudaDeviceDisablePeerAccess(0); // dev 0 is no longer a peer of dev 1
其他并发场景
主机/设备执行并发
1
2 | Kernel<<<b, t>>>(...); // 这个内核执行可以与
cpuFunction(...); // 这段主机代码重叠
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
