cuda-notes

杂记

1.可以在kernel 里设置同步点进行数据间的同步，__syncthreads()通常用于调整在相同块之间的线程通信 2.一个批处理中每一个多处理器可以处理多少个块，取决于每个线程中分配了多少个寄存器和已知内核中每 个时钟需要多少的共享内存

3.warp是硬件层面中SM对应执行线程的单位。

线程束Wrap是GPU的基本执行单元，目前cuda的warp的大小为32。

同在一个warp的线程执行同一指令。warp 本质上是一组被同时调用的线程组，符合SIMD（或叫 SIMT 单执行多线程）并行模型，所有线程中都执行相同的指令，每一条warp 包含数量相同的线程；warp周期性切换执行；关键要理解相同指令的执行，也就意味着如果出现不同的执行分支，同一 warp 中的一部分线程会阻塞等待一另部分执行完才能执行后面的相同指令

4.线程块是一个逻辑上的概念，其大小可以设置，其中的线程被分配一块共享内存，其线程数往往大于warp 大小（常规cuda 设计下，warp size为32），也就意味着一个块中的会划分出多个warp。线程块中的线程数量曾经受架构限制，每个线程块最多只有512个线程，但从2019年7月开始，线程块可以最多包含1024个线程。

wiki 线程块

与线程块对应的硬件是SM（stream multiprocessor）, 不同的 SM 之间无法进行同步，因此一个block 只能分配到一个SM 中，相反的，一个 SM 却可以承载多个block 设置block size 和 grid size

5. 一个块内的 warp 次序是未定义的，但可以通过协调全局或共享内存的读取进行同步；在一个栅格块内的块次序是未定义的，并且不存在同步机制

6.内存模型 一条执行在设备上的线程，只允许通过如下的内存空间使用设备的DRAM 和On-Chip 内存，如图2-2 所 示：
读写每条线程的寄存器，
读写每条线程的本地内存，
读写每个块的共享内存，
读写每个栅格的全局内存，
只读每个栅格的常量内存，
只读每个栅格的纹理内存。

全局内存在设备的 DRAM 上；常量内存和纹理内存都是只读内存 全局内存不会被缓存；常量内存和纹理内存中的数据会被缓存；缓存是无法被直接访问的，cuda 中的缓存数据来源是常量内存和纹理内存

7.循环展开的： #pragma unroll 默认下，编译器为已知的行程计数展开小型循环。 #pragma unroll 5 For (int i = 0; i < n; ++i) 循环将展开5 次。 如果#pragma unroll 后面没有附值，当行程计数（n）为常数时，循环完全展开，否则不会展 开。

8.-use_fast_math编译选项，将替换math 库中的标准标本为低准确性而更快速的版本

9.page-locked memory： 在cpu 内存中分配的内存，不会被换出，gpu 使用时通过PCIE 进行通信 10.unified memory(managed memory)： 可以翻译为托管内存，将分配的驻留位置移动到需要它的处理器，当gpu 需要访问时就会把该块内存数据移动到gpu 内存上，当 cpu 需要访问时就把这块内存数据移动到cpu 内存上。

11.以下操作在不同的流中也不能进行并发：page-locked 主机内存分配，设备内存的分配，设备内存的 设置，设备到设备的内存拷贝，或它们之间的事件纪录。 12.数据拷贝和数值计算、两个方向的拷贝（gpu->cpu 和 cpu->gpu，两个gpu->cpu 或者一个 gpu->cpu、一个cpu->gpu这种）可以同时进行。 13.cudaThreadSynchronize调用后可以确保前面所有的流都已完成。

12.设置CUDA_LAUNCH_BLOCKING 环境变量为1，强制cuda同步运行

13.为保证编译器生成最低数量的指令，对于结构体大于16 字节的，应该用__align__ (16)定义,例如:

struct __align(16)__ {
  float a;
  float b;
  float c
  float d
  float e;
};
被编译成为二个128-bit 加载指令而不是五个32-bit 加载指令。 

14.bank 冲突分两种，一种是共享内存bank 冲突，共享内存在物理上被分为 32 个（刚好等于一个线程束中的线程数目，即内建变量 warpSize 的值）同样宽度的、能被同时访问的内存 bank，当同一个 warp 中的不同线程访问同一个 bank 中的不同的地址时（访问同一个地址则会发生广播），就会发生 bank 冲突；另一种是寄存器 bank 冲突，每个 SM 有独立的 register file，每个reigister file 被分为多个 bank，寄存器编号 id%4 即为该寄存器所属的 bank，如 R0 属于 Bank 0，R5 属于 Bank 1），当一条指令所需的源寄存器有两个以上来自于同一 bank则会产生 conflict，比如FFMA R1, R0, R4, R1 这样的指令就回产生寄存器 Bank conflict（同同访问了 R0 和 R4，同属于 bank 0 的寄存器）

http://www.zh0ngtian.tech/posts/96744e8c.html

15. pitch内存和linear 内存的区别 什么情况下使用cuMemAllocPitch，cuMemAllocPitch分配的内存可以和cudaMalloc分配的内存进行互相替代吗 ：

cuMemAllocPitch 主要在以下情况下使用： 处理2D数组/图像数据： 当需要处理2D数据（如图像）时，cuMemAllocPitch会自动处理内存对齐问题 它会在每行末尾添加额外的填充字节，以确保每行的起始地址都能够正确对齐，这有助于提高内存访问性能

优化内存访问： 通过确保每行数据都按照适当的边界对齐，可以实现更高效的内存访问 对于某些硬件架构，这种对齐可以显著提升性能

关于是否可以与cudaMalloc互相替代： 不能直接替代，原因如下： cuMemAllocPitch分配的内存包含额外的填充字节，实际占用的内存空间大于原始数据大小 访问cuMemAllocPitch分配的内存需要使用返回的pitch值来计算正确的内存偏移 数据布局不同，cudaMalloc分配的是连续的内存空间，而cuMemAllocPitch包含了行间填充

示例代码

// 使用cudaMalloc
float* data;
int width = 1024;
int height = 768;
cudaMalloc((void**)&data, width * height * sizeof(float));

// 使用cuMemAllocPitch
float* data;
size_t pitch;
int width = 1024;
int height = 768;
cuMemAllocPitch((CUdeviceptr*)&data, &pitch, width * sizeof(float), height, 16);

// cudaMalloc方式访问
__global__ void kernel1(float* data, int width) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int idx = y * width + x;
    float value = data[idx];
}

// cuMemAllocPitch方式访问
__global__ void kernel2(float* data, size_t pitch, int width) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    float* row = (float*)((char*)data + y * pitch);
    float value = row[x];
}

总结： 如果您主要处理2D数据（特别是图像处理），并且关注内存访问性能，建议使用cuMemAllocPitch 如果是处理一维数组或者不需要特别考虑内存对齐，使用cudaMalloc会更简单直接 两者不能直接互相替代，因为内存布局和访问方式都不同

什么情况下进行互相替代使用呢

cuMemAllocPitch 和 cudaMalloc 在以下情况下可以考虑互相替代使用： 一维数组场景： pitch 恰好等于宽度的情况：

// 如果通过 cuMemAllocPitch 得到的 pitch 值恰好等于 width * sizeof(float)
float* data;
size_t pitch;
int width = 128;  // 假设这个宽度正好满足对齐要求, 宽度为1920 时，pitch 为2048，此时和 malloc 分配的内存不一样
int height = 768;
cuMemAllocPitch((CUdeviceptr*)&data, &pitch, width * sizeof(float), height, 16);

// 如果 pitch == width * sizeof(float)，则可以用 cudaMalloc 替代
float* data;
cudaMalloc((void**)&data, width * height * sizeof(float));

数据拷贝后重新组织的场景：

// 原始数据使用 cuMemAllocPitch 分配
float* src_data;
size_t src_pitch;
cuMemAllocPitch((CUdeviceptr*)&src_data, &src_pitch, width * sizeof(float), height, 16);

// 可以将数据拷贝到 cudaMalloc 分配的连续内存中
float* dst_data;
cudaMalloc((void**)&dst_data, width * height * sizeof(float));

// 拷贝时重新组织数据
for(int y = 0; y < height; y++) {
    cudaMemcpy(dst_data + y * width, 
               (char*)src_data + y * src_pitch, 
               width * sizeof(float), 
               cudaMemcpyDeviceToDevice);
}