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...
constexpr, group_size_m: tl.constexpr, fp8_fast_accum: tl.constexpr,): 这个函数的输入包括: a_ptr, b_ptr, c_ptr:分别是矩阵 A、B 和 C 的指针(分别表示输入矩阵和输出矩阵)。 scales_ptr, zeros_ptr:分别是量化时使用的比例因子和零值,用于恢复量化矩阵的精度。 stride_* 参数:用于访问矩阵元素...
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 ...
BLOCK_SIZE_N:tl.constexpr,# BLOCK_SIZE_K:tl.constexpr,# GROUP_SIZE_M:tl.constexpr,#):pid=tl.program_id(axis=0)num_pid_m=tl.cdiv(M,BLOCK_SIZE_M)num_pid_n=tl.cdiv(N,BLOCK_SIZE_N)num_pid_in_group=GROUP_SIZE_M*num_pid_n ...
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),其中因为我们预先认为 ...