| 244 | |
| 245 | |
| 246 | def matmul248(input, qweight, scales, qzeros, g_idx, bits, maxq): |
| 247 | output = torch.empty((input.shape[0], qweight.shape[1]), device='cuda', dtype=torch.float16) |
| 248 | grid = lambda META: ( |
| 249 | triton.cdiv(input.shape[0], META['BLOCK_SIZE_M']) * triton.cdiv(qweight.shape[1], META['BLOCK_SIZE_N']),) |
| 250 | matmul_248_kernel[grid](input, qweight, output, |
| 251 | scales, qzeros, g_idx, |
| 252 | input.shape[0], qweight.shape[1], input.shape[1], bits, maxq, |
| 253 | input.stride(0), input.stride(1), |
| 254 | qweight.stride(0), qweight.stride(1), |
| 255 | output.stride(0), output.stride(1), |
| 256 | scales.stride(0), qzeros.stride(0)) |
| 257 | return output |
| 258 | |
| 259 | |
| 260 | def transpose_matmul248(input, qweight, scales, qzeros, g_idx, bits, maxq): |