GROUP_SIZE_M = 64 if N <= 8192: GROUP_SIZE_M = 96 if N <= 4096: GROUP_SIZE_M = 128 if N <= 1024: GROUP_SIZE_M = 256 # allocate output # 分配输出 locks = torch.zeros(2 * GROUP_SIZE_M, dtype=torch.int32, device=w.device) _dw = torch.zeros((GROUP_SIZE_M, N), dty...
first_pid_m = group_id * GROUP_SIZE_M # If `num_pid_m` isn't divisible by `GROUP_SIZE_M`, the last group is smaller # 如果 `num_pid_m` 不能被 `GROUP_SIZE_M` 整除,最后一组会比较小 group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) # *Within groups*, progra...
triton.Config({'BLOCK_SIZE_M':128,'BLOCK_SIZE_N':128,'BLOCK_SIZE_K':32,'GROUP_SIZE_M':8},num_stages=4,num_warps=4), triton.Config({'BLOCK_SIZE_M':128,'BLOCK_SIZE_N':64,'BLOCK_SIZE_K':32,'GROUP_SIZE_M':8},num_stages=4,num_warps=4), triton.Config({'BLOCK_SIZE_M':6...
first_pid_m = group_id * GROUP_SIZE_M # If `num_pid_m` isn't divisible by `GROUP_SIZE_M`, the last group is smaller # 如果 `num_pid_m` 不能被 `GROUP_SIZE_M` 整除,最后一组会比较小 group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) # *Within groups*, progra...
group_size_m=min(num_pid_m-first_pid_m,GROUP_SIZE_M) pid_m=first_pid_m+(pid%group_size_m) pid_n=(pid%num_pid_in_group)//group_size_m #--- #CreatepointersforthefirstblocksofAandB. #WewilladvancethispointeraswemoveintheKdirection...
设输入行数 M=4 和 GROUP_SIZE_M=2,以下是 ∇w 的并行归约策略图示(为简洁起见,省略 ∇b): 在第一阶段,同色的 X 行共享同一个缓冲区,因此使用 lock 以确保一次只有一个内核实例写入缓冲区。在第二阶段,这些缓冲区会进一步归约以计算最终的 ∇w 和 ∇b。在以下实现中,第一阶段由函数 _layer_...
GROUP_SIZE_M:tl.constexpr,BLOCK_SIZE_N:tl.constexpr):#Map the program id to the elements of X,DX,and DY it should compute.# 映射程序 id 到对应计算的 X,DX,DY row=tl.program_id(0)cols=tl.arange(0,BLOCK_SIZE_N)mask=cols<N ...
group_a_ptrs,group_b_ptrs,group_c_ptrs,#devicetensor of gemm sizes.its shape is[group_size,3]# 设备张量的 GEMM(General Matrix Multiply)大小。其形状为[group_size,3]#dim0is group_size,dim1is the values of<M,N,K>of each gemm# 第0维是 group_size,第1维是每个 GEMM 的<M,N,K>...
group_a_ptrs,group_b_ptrs,group_c_ptrs,# device tensorofgemm sizes.its shape is[group_size,3]# 设备张量的GEMM(General Matrix Multiply)大小。其形状为[group_size,3]# dim0is group_size,dim1is the valuesof<M,N,K>ofeach gemm
在上面的示例A(9, 9) @ B(9, 9) = C(9, 9)中,因为BLOCK_SIZE_M 、BLOCK_SIZE_N以及BLOCK_SIZE_K的设置都是3,并且实际上在上面的例子中GROUP_SIZE_M的设置是1,实在太特殊,所以我们再举个数值上更能体现分组的例子。 A(2048, 2048) @ B(2048, 2048) = C(2048, 2048),其中因为我们预先认为 ...