目录
前言
我们所有的优化方式是从顶层的，从之前的展开也好，再到后来的内存方面，这一章我们将详细的谈一下核函数，具体来说不是之前的那种展开，让SM达到最大利用率，而是数据维度方面，让带宽达到最大
GDDR6
前面没有详细谈，这里详谈一下，不然老是作为缺失的一角
GDDR6是一种硬件技术标准，它和显存颗粒的关系就好像usb协议和usb接口的关系
显存颗粒之间的通信就是靠GDDR6
它是显存颗粒的“制造规范”：GDDR6（Graphics Double Data Rate 6）是由行业标准组织 JEDEC 制定的第六代图形用双倍速率同步动态随机存储器标准。它严格规定了显存颗粒的物理尺寸、电气特性、数据传输速率、电压等物理层参数。这就像 USB-IF 组织制定了 USB4 标准一样，是硬件设计蓝图。
它体现在物理实体上：显卡上围绕在 GPU 核心周围的黑色小方块就是 GDDR6 显存颗粒。这些颗粒由三星、美光、海力士等厂商按照 GDDR6 标准制造。同时，GPU 核心内部也集成了支持 GDDR6 协议的显存控制器，负责与这些颗粒通信。
它由硬件电路实现，而非软件驱动：GDDR6 的核心技术，如双倍数据速率、突发传输、PAM4 信号调制等，都是由 GPU 芯片和显存颗粒内部的物理电路直接完成的，不需要任何软件驱动介入。驱动程序只负责初始化显存控制器，配置时序参数，一旦配置完成，所有数据搬运都是纯硬件行为，和我们之前讨论的 LSU、L2 缓存一样，都是硬件在自动工作。
GDDR6 决定了“每条数据线的速度”： RTX 4060 的 GDDR6 单线速率是 17 Gbps。
显存位宽决定了“有多少条数据线”： RTX 4060 的位宽是 128 位（128 条线）。
两者共同决定显存带宽：272 GB/s = 17 GHz × 128 bit / 8。
┌──────────────┐ │ GPU 核心 │ │ 显存控制器 │ │ 128 位总位宽 │ └──┬──┬──┬──┬──┘ │ │ │ │ ← 每个颗粒独占 32 位数据线 ┌──┘ │ │ └──┐ │ │ │ │ 颗粒0 颗粒1 颗粒2 颗粒3 (32位) (32位) (32位) (32位)如果按照我的4060 8GB显存来看，一个颗粒就是2GB，每一颗GDDR6显存颗粒就是独立的小仓库，然后拉出去32条数据线，链接内部的GPU，那就是32*4=128位
这四个就是显存颗粒
显存（GDDR6）在物理上不是一整块连续的存储 而是被分成多个独立的存储单元，层级如下： Channel（通道） ↓ Memory Partition（内存分区） ↓ Bank Group（bank组） ↓ Bank（bank） ↓ Row / Column（行列） 材料仓库（全局内存） ├── 仓库大厅A（Channel 0） │ ├── 库房1（Bank Group 0） │ │ ├── 小格子1（Bank 0） │ │ ├── 小格子2（Bank 1） │ │ ├── 小格子3（Bank 2） │ │ └── 小格子4（Bank 3） │ └── 库房2（Bank Group 1） │ └── ... └── 仓库大厅B（Channel 1） └── ...注意取货的时候同一间的bank是打开一行，也就是Row，如果你的数据同时在一个bank里面的不同Row，只能串行访问，也就是会bank conflict，它只能同一时刻打开一行，然后再根据column去定位某个字节，就会导致你Row的时候，其他的不能来，必须放回去之后下一个请求才能来
4060是4*16=64个独立的bank根据以上的只是，我们需要放数据放在同一个bank的同一行，如果你是同一个bank的不同行，是需要先关闭一行然后再打开一行的，如果你是不同的bank，那本质是独立的互不冲突
所以我们的数据如果是同一个bank的同一行，那读取这一行的时候效率会非常高
为什么这样分？？？
并行并行，这样能够调度起来的话，能够极大的提高并行性
所有的优化方向无非三点
1：增大带宽，换言之就是数据线上跑的车，运输物料的车要足够的多，要一直再跑，不能阻塞停下来
2：让工厂里面的工人永不停息，就是cuda核心的占有率高，全负荷运行
3：如果12满足了，那只剩下算法了，也就是你设计的算法效率是不是高效的
内存延迟和内存带宽
内存延迟关注的是“等多久”，是指从发出一个内存访问请求开始，到数据真正准备就绪（比如加载到寄存器）为止，总共花费的时钟周期数。它衡量的是数据访问的响应时间。
内存带宽关注的是“搬多快”，是指单位时间内（通常是每秒）内存系统能够传输的数据总量，比如 GB/s。它衡量的是数据访问的吞吐速率。
理论带宽就是之前提到的硬件272GB/s
概念 衡量什么 单位 对应 ncu 指标 比喻 有效带宽 程序真正需要的数据传输速率 GB/s 有效需求字节数 / 时间 工厂真正加工的有效原料的运送速度 内存吞吐量 硬件实际搬运的总数据速率 GB/s dram__bytes_read.sum.per_second传送带实际运送的总物料速度（含废料） 计算吞吐量 计算核心的算力 GFlops sm__throughput.avg.pct_of_peak_sustained_elapsed工厂加工原料的速度
通过以上的，我们完全可以找到优化思路，比如
先看有效带宽是否远低于理论带宽 → 如果是，说明访存效率低，瓶颈在数据搬运。
再看内存吞吐量是否接近理论带宽 → 如果是，说明硬件已经在满负荷搬东西，但搬了很多无用数据。此时优化的重点是提高有效数据占比（合并访问、对齐）。
如果内存吞吐量本身就很低 → 说明硬件闲着，需要增加并发度（提高占用率、增加Warp数量）。
如果有效带宽已经接近理论带宽，但程序还是慢 → 瓶颈不在带宽，而在计算核心。此时要看计算吞吐量，优化计算部分。
矩阵转置问题
过去我们常常研究的是如何最大化线程束来隐藏延迟+对齐合并
但是如果内核写的不好，那无疑是在拖拉机上面进行动力学优化，我们做的就是改造拖拉机本身，为此引入一个矩阵转置问题的背景
矩阵转置是所有学理工科的线性代数里面的某一个小的知识点
简单来说就是把行变成列
__global__ void transpose_naive(float *odata, const float *idata, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { odata[x * height + y] = idata[y * width + x]; } }那么内存当中实际的布局就是这样
通过观察可以发现，输出odata不是合并的，输入idata是合并的
那么有什么办法让两者都能合并呢？使得访存的效率得到大大提升
列读取
上面的是行读取，也就是可以合并，但是如果采用列读取呢？
如果一开始第一列读取了之后，会把相邻的给读取进来，然后由于L1缓存存在，就是把相邻的放在缓存里面，导致第二列第三列的读取就不需要去全局内存了，这样的速度会不会更快？并且还有写分配的问题，如果写是交叉的，那会导致大量的写分配
计算上下限
很明显，行会很强，所以按照
行读取+行存储=上限
列读取+列存储=下限
我们可以写个函数计算一下上下限，注意我们仅仅是模拟访存模式，而不是转置
暂时没有写cpu端，因为我们的上下限计算的时候不是转置，只是单纯的访问形式去计算
#include "../freshman.hpp" #include <cuda_runtime.h> //行存储+读取 __global__ void row_bound(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height){ int idx = y * width + x; odata[idx] = idata[idx]; // 读和写都是同一行连续地址 } } // 列存储+读取 __global__ void column_bound(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height){ int idx = x * height + y; // 注意：这里把width换成了height odata[idx] = idata[idx]; } } void cpuMatrix(float* odata,float* idata,int width,int height){ } int main(){ int device=0; cudaSetDevice(device); int tran_kernel=1;//给后续的switch语句控制的，能够挑选运行的核函数 //数据量 int width=1<<12; int height=1<<12; int nElem=width*height; int nSize=sizeof(float)*nElem; //开辟cpu内存，初始化数据 float*a_h,*b_h; //把矩阵b转换为a float* res_h; a_h=(float*)malloc(nSize); b_h=(float*)malloc(nSize); res_h=(float*)malloc(nSize); initialData(b_h,nElem); //开辟gpu内存，把b_h的数据拷贝过去 float* b_d; float* res_d; CHECK(cudaMalloc((void**)&b_d,nSize)); CHECK(cudaMalloc((void**)&res_d,nSize)); CHECK(cudaMemcpy(b_d,b_h,nSize,cudaMemcpyHostToDevice)); //CPU端 // double start=efficiency(); // cpuMatrix(a_h,b_h,width,height); // double end=efficiency(); // GPU端 int dimx = 32; int dimy = 32; dim3 block(dimx, dimy); dim3 grid((width + block.x - 1) / block.x, (height + block.y - 1) / block.y); double startGpu=efficiency(); switch (tran_kernel){ case 0: row_bound<<<grid,block>>>(res_d,b_d,width,height); break; case 1: column_bound<<<grid,block>>>(res_d,b_d,width,height); break; } cudaDeviceSynchronize(); double endGpu=efficiency(); CHECK(cudaMemcpy(res_h,res_d,nSize,cudaMemcpyDeviceToHost)); //check(res_h,res_d); //std::cout<<"CPU elapsed: "<<end-start<<std::endl; std::cout<<"GPU elapsed: "<<endGpu-startGpu<<std::endl; free(a_h); free(b_h); free(res_h); cudaFree(b_d); cudaFree(res_d); return 0; }一共运行6遍，去除冷启动，实际5遍，然后取平均值，所以我们的交叉访问的性能应该介于两者之间，这将作为我们后续的指导，不断的把优化接近上限
指标 数值 (秒) 说明 上限（行读行写） 0.001614 秒 (平均) 最优访存模式，零交叉访问，带宽利用率最高 下限（列读列写） 0.002584 秒 (平均) 最差访存模式，全交叉访问，带宽浪费严重 性能差距 约 1.60 倍 上限比下限快 60%，这就是交叉访问导致的性能损失上限
套入有效带宽的计算公式即可得到
版本 平均时间 有效带宽 上限（行读行写） 0.001614 秒 83.15 GB/s 下限（列读列写） 0.002584 秒 51.94 GB/s
朴素转置
分别采取行读取和列读取的方式进行转置，这也是最容易想到的办法
//朴素转置：行读取 __global__ void transform_row(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height){ odata[x * height + y] = idata[y * width + x]; } } //朴素转置：列读取 __global__ void transform_col(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height){ odata[y * width + x] = idata[x * height + y]; } }
版本 实测时间（取稳定值） 相当于哪个操作？ 控制函数为：2（行读取的转置） 约 0.00230~0.00255 秒 合并读 + 交叉写 控制函数为：3（列读取的转置） 约 0.00155~0.00198 秒 交叉读 + 合并写
版本 平均时间（取稳定值） 估算有效带宽 行读取的转置（合并读+交叉写） 约 0.00240 秒 0.1342 / 0.00240 ≈ 55.9 GB/s 列读取的转置（交叉读+合并写） 约 0.00170 秒 0.1342 / 0.00170 ≈ 78.9 GB/s 
对于读取来说：瓶颈在于是否合并读取，行读取时是合并的，但是列不是，所以行读取时原料能够快速的送达到cuda核心进行计算，但是列读取有个优势，就是可能会预读取，就是把周围的也加载到l1l2缓存，下次就没必要去显存，所以也会提升效率，并且如果读受阻了，还可以利用warp的并行去掩盖延迟
对于写入来说：瓶颈在于是否合并对齐，列读取时后续的写入是可以合并的，但是行读取后续的写入不是合并的，但是注意还有个写分配问题，对于行读取后续不合并的写入来说写分配耗时很大
还需要注意 cuda计算是跟原料有关系，所以原料送的越快，那cuda就可以很快开始计算，但是后续的写入只要发出请求，LSU受理之后跟cuda就没关系了，写到缓存之后的写分配之类的消耗是跟cuda没关系的，cuda可以去计算别的
综上，行列读取都各有优势，但以博主微薄的知识来看，应该是写分配消耗太多的时间了交叉读（读不合并）的代价： 多从DRAM搬了数据（163 GB/s vs 82 GB/s） 但GPU有大量warp，当一个warp在等数据时 调度器切换到其他warp继续执行 → 延迟被并发的warp掩盖掉了 → 虽然搬运量大，但时间不一定长 而且这里面虽然第一轮的读不合并，但是数据又缓存到L1L2缓存当中，这一点收益还是有的 交叉写（写不合并）的代价： 触发大量写分配 → 每次写都要先读一个cache line进来 这些"写分配导致的读"是真实的读操作 同样会阻塞线程（读必须等）（这里的意思是因为写分配的读导致其他需要进入cuda核心的数据被阻塞） 相当于"把写不合并的代价转化成了读不合并的代价" 但这些额外的读是在关键路径上的，无法被掩盖 → 性能损失更严重 transform_row（合并读+交叉写）： 读高效（sector/request=4） 但交叉写触发大量写分配→产生隐性读阻塞→性能受限 → 耗时0.0024s transform_col（交叉读+合并写）： 读低效（sector/request=32）→搬运大量无用数据（163 GB/s） 但合并写减少写分配→写操作fire and forget不阻塞 交叉读的延迟被大量并发warp掩盖 → 耗时0.0017s，反而更快 核心结论： 写操作的合并比读操作的合并更重要 因为写不合并会通过写分配转化为读阻塞 而读不合并的延迟可以被warp并发掩盖经过以上，如果只能选择一种，尽量选择写合并，写合并比读合并好一点
展开转置
//展开转置：行读取 __global__ void transform_unroll_row(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x*4 + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { odata[x * height + y] = idata[y * width + x]; if (x + blockDim.x < width) odata[x * height + y + height*blockDim.x] = idata[y * width + x + blockDim.x]; if (x + blockDim.x*2 < width) odata[x * height + y + height*blockDim.x*2] = idata[y * width + x + blockDim.x*2]; if (x + blockDim.x*3 < width) odata[x * height + y + height*blockDim.x*3] = idata[y * width + x + blockDim.x*3]; } } //展开转置：列读取 __global__ void transform_unroll_col(float *odata, float *idata, int width, int height){ int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y * 4 + threadIdx.y; if (x < width && y < height) { odata[x * height + y] = idata[y * width + x]; if (y + blockDim.y < height) odata[x * height + y + blockDim.y] = idata[(y + blockDim.y) * width + x]; if (y + 2 * blockDim.y < height) odata[x * height + y + 2 * blockDim.y] = idata[(y + 2 * blockDim.y) * width + x]; if (y + 3 * blockDim.y < height) odata[x * height + y + 3 * blockDim.y] = idata[(y + 3 * blockDim.y) * width + x]; } }行读
- CPU 平均耗时：0.22906 s
- GPU 平均耗时：0.007338 s
列读
- CPU 平均耗时：0.202909 s
- GPU 平均耗时：0.007457 s
实验的波动非常大，有时候达到0.01的量级，有时候达到0.001的量级，五次取平均，不知道是因为wsl的关系还是因为啥（本身的核函数设计的也不合理）
对于展开来说：收益来源于三点，可以联系之前的加法
- 减少循环控制指令开销；
- 提升指令级并行（ILP），隐藏指令延迟；
- 配合共享内存减少
__syncthreads同步次数。但在这里没有循环，展开实际的收益不大，ILP是针对那些上下有依赖的计算密集型提升很大，这里主要是访存延迟问题，转置不像加法依赖上一轮，如果依赖上一轮，你可以直接在一轮里面多算，然后指令并行，就可以隐藏指令延迟，收益非常大，但是这里的收益可以算作微乎其微，和之前都没啥变化，第三点这里也没有
而且block的数量变少，导致能够隐藏延迟的能力变差
实际的问题还是在访存密集型而不是计算密集型，每次访存都是几百的cycle，我们要解决的是访存的问题
对角转置
这里是针对dram的访存进行优化的
GDDR6 单颗颗粒物理封装对外是 32bit 总线，但内部拆成2 个完全独立 16bit 片内通道（Channel A / Channel B），各有独立地址、命令、数据线、Bank 组，能并行收发请求
① GPU 全局 Channel：GPU 内存控制器分出的通道，一颗颗粒对应一个；
② 颗粒内部子通道：单颗 DRAM 芯片内部拆分的 2 个 16bit 通道，属于颗粒内部硬件。
单颗GDDR6颗粒 ├─ 片内通道A（16bit） │ └─ Bank Group 0~1（共8个Bank） ├─ 片内通道B（16bit） │ └─ Bank Group 2~3（共8个Bank） └─ 总计：16个Bank，分为4个Bank Group（每组4个Bank）内存分区就是一个显存颗粒一个全局channel，然后一个显存颗粒内部再细分两个子channel
一个bank是由行列组成的二维存储阵列，之前提到过如果你访问同一个bank的不同行，必须串行，也就是必须关闭一行在打开零一行
- Row（行） 一次 ACTIVATE（激活）命令会把一整行数据全部加载到该行缓冲区（Row Buffer，相当于 Bank 内部高速缓存）；同一 Bank 频繁切换行会产生巨大延迟（行冲突）。
- Column（列） 行激活后，通过列地址读取 / 写入行缓冲区里指定位置的数据；GDDR6 单次突发 16 个字，一次传输 32 字节。
如果是同一组的不同bank需要两次激活，可以并行，但是性能肯定弱于不同组的bank，不同组只需要一次激活，所以性能最大化就是不同组
读取一行到缓冲区之后，是到列读取，列读取某一段32字节（突发传输），那两个channel，就是64字节，64*4=256字节，注意这不是一次性输出，我们的数据位只有128bit，这需要多次的输出，这两不是同一个层次的，不要弄混
所以综上，我们的目的就是访问不同的bank地址是连续的，如果我们的线程访问的地址也是连续的，那可能都会打在同一个bank中
关于线程访问的地址都是我们人为规定的，你当然可以让线程1访问arr[3]，线程2访问arr[10]，完全打乱，我们之前按照线程1访问arr[1]，线程2访问arr[2]这样都是我们认为规定的，所以有没有一种办法，尽量的打乱，让其访问不同的bank？？？
对角转置的目的就是：让相邻线程块访问的全局内存地址分散到不同的DRAM Bank上，让所有Bank都能并行工作，最大化带宽利用率。（注意这里是线程块哦，这是一个宏观层面，也就是发出请求，我们看请求是否都在同一个bank，我们要避免，比如线程块1发的请求要去channel0，而线程块2也去channel0，那不就冲突了，我们要避免就是这种情况）
f(x, y) = (m, n)我们只需要找到一种函数，原来线程块的坐标是（x,y），我们映射到新的(m,n)上面，然后按照新的坐标去访问内存
原理：RTX 4060显存： - 4个channel（通道） - 每个channel有多个bank - 地址按照某种规则分配到不同channel/bank 关键规则：相邻的256字节（或512字节）地址被分配到不同的channel 这叫做地址交织（address interleaving） 例子（简化）： 地址0~255 → channel 0 地址256~511 → channel 1 地址512~767 → channel 2 地址768~1023 → channel 3 地址1024~1279→ channel 0（循环） ...我们的程序采用的是虚拟地址，block_idx=虚拟地址/256，得到一个数字，然后再%4，也就是这256字节分配到哪个channel，这个就叫地址交织，为了让地址访问的时候能够分配在不同的channel，提高并行性
那么原理就是让访问的地址分在不同的channel，但是我们之前的交错访问可能导致访问同一个channel
比如：之前的朴素行转置
矩阵4096×4096，block=32×32 block(0,0)：读A[0~31][0~31] → 写B[0~31][0~31] block(1,0)：读A[0~31][32~63] → 写B[32~63][0~31] block(2,0)：读A[0~31][64~95] → 写B[64~95][0~31] ... block(127,0)：读A[0~31][4064~4095] → 写B[4064~4095][0~31] 注意看写操作： B[0~31][0~31] 的起始地址：0 B[32~63][0~31] 的起始地址：32*4096*4 = 524288字节 B[64~95][0~31] 的起始地址：64*4096*4 = 1048576字节 ... 所有block同时写B的第0~31列 这些写操作的地址： B[0][0] 地址：0字节 B[32][0] 地址：524288字节 = 512K字节 B[64][0] 地址：1048576字节 = 1024K字节 B[96][0] 地址：1572864字节 = 1536K字节 524288 / 256 = 2048个channel步长 2048 % 4 = 0 → 全部落在channel 0！ 所有block的写操作全部竞争同一个channel → channel 0严重拥堵 → channel 1、2、3完全空闲 → 4个channel只用了1个，带宽利用率25%朴素版（列顺序）： 对角线版： ┌─────────────────┐ ┌─────────────────┐ │ 0 1 2 3 ... │ │ 0 1 2 3 ... │ │ 0 1 2 3 ... │ │ 1 2 3 0 ... │ │ 0 1 2 3 ... │ │ 2 3 0 1 ... │ │ 0 1 2 3 ... │ │ 3 0 1 2 ... │ └─────────────────┘ └─────────────────┘按照block.x递增来调度的化，处于第0行的block，他们在写的时候就会导致在同一个channel
解决方案：
int block_x = (blockIdx.x + blockIdx.y) % gridDim.x; // 列方向block索引（关键） int block_y = blockIdx.x; // 行方向block索引朴素版（列顺序）： 对角线版： ┌─────────────────┐ ┌─────────────────┐ │ 0 1 2 3 ... │ │ 0 1 2 3 ... │ │ 0 1 2 3 ... │ │ 1 2 3 0 ... │ │ 0 1 2 3 ... │ │ 2 3 0 1 ... │ │ 0 1 2 3 ... │ │ 3 0 1 2 ... │ └─────────────────┘ └─────────────────┘这样就会变成右边的对角线版本
简单来理解就是线程块还是原来的线程块，但是处理的数据位置变了，以写操作来看，也就是线程块(0,1)原来处理数据块(1,0)的数据，但是现在线程块(0,1)来处理数据块(1,1)的数据
//对角转置 __global__ void transform_diagonal_row(float *odata, float *idata, int width, int height) { int block_y=blockIdx.x; int block_x=(blockIdx.x+blockIdx.y)%gridDim.x; int ix=threadIdx.x+blockDim.x*block_x; int iy=threadIdx.y+blockDim.y*block_y; int idx_row=ix+iy*width; int idx_col=ix*height+iy; if (ix<width && iy<height) { odata[idx_col]=idata[idx_row]; } } __global__ void transform_diagonal_col(float *odata, float *idata, int width, int height) { int block_y=blockIdx.x; int block_x=(blockIdx.x+blockIdx.y)%gridDim.x; int ix=threadIdx.x+blockDim.x*block_x; int iy=threadIdx.y+blockDim.y*block_y; int idx_row=ix+iy*width; int idx_col=ix*height+iy; if (ix<width && iy<height) { odata[idx_row]=idata[idx_col]; } }可以从代码看出，我们改变的仅仅是odata[]和idata[]的访问，我们压根没有动原来的线程块
换言之就是我们让线程块的idx不是一一对应数据块的idx，而是让线程块访问不同的数据块
另外博主测试的时候发现，环境极其不稳定，发现展开的实验数据又有变化，所以重新测了一下同一时刻下的数据，发现性能：7>6和5>4
对角坐标变换 (函数6) 0.003443 秒 对角坐标变换 (函数7) 0.002588 秒
基本可以得出来一个结论，就是行读列写的一大准则，这个列写基本都是最快的
注意一点：调度队列规则
GPU 硬件有全局的线程块分发器（Distributor）：
- 所有 grid 里的 thread block，按
blockIdx.x → blockIdx.y → blockIdx.z顺序进入待调度队列； 逻辑上队列是有序排队的，先创建的 block 排在前面。- 但是分发逻辑是抢占式空闲调度： 只要任意一个 SM 内部有空余资源（寄存器、共享内存、warp 槽位），分发器就把队列最前面的 block 派送到这个空闲 SM。
那么我们的block究竟有没有像我们之前说的那样已经按照blockIdx.x去调度了？有没有变化？这也是一个值得思考的问题，如果本身被打断了的话，那会不会本身的访问已经打在了不同的channel上，我们的对角转置又恰好回来了访问到同一个channel上面，似乎是个概率问题？还有博主后台开了很多的软件加上wsl之类的，可能会有很大的波动，这可能也会影响调度，博主知识有限，只能分析到这里
总结
以上的所有实验，都是在说明带宽这一路径，也就是我们前面分析的三条优化方向
1：带宽
2：算法
3：cuda占有率
这里我们是不断的改进核函数访存效率，提升带宽从而提升效率
遗憾的是博主跟着谭升大佬学习，可能是博客过老讲解的是老的架构，跟新架构的出入，或者是环境的不同，所以有些实验是有出入的，但是整体的知识框架是对的，也欢迎大家指正
374
扫一扫
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
