| 258 | |
| 259 | |
| 260 | def transpose_matmul248(input, qweight, scales, qzeros, g_idx, bits, maxq): |
| 261 | output_dim = (qweight.shape[0] * 32) // bits |
| 262 | output = torch.empty((input.shape[0], output_dim), device='cuda', dtype=torch.float16) |
| 263 | grid = lambda META: ( |
| 264 | triton.cdiv(input.shape[0], META['BLOCK_SIZE_M']) * triton.cdiv(output_dim, META['BLOCK_SIZE_K']),) |
| 265 | transpose_matmul_248_kernel[grid](input, qweight, output, |
| 266 | scales, qzeros, g_idx, |
| 267 | input.shape[0], qweight.shape[1], output_dim, bits, maxq, |
| 268 | input.stride(0), input.stride(1), |
| 269 | qweight.stride(0), qweight.stride(1), |
| 270 | output.stride(0), output.stride(1), |
| 271 | scales.stride(0), qzeros.stride(0)) |
| 272 | return output |
| 273 | |
| 274 | |
| 275 | class QuantLinearFunction(torch.autograd.Function): |