You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "masahi (via GitHub)" <gi...@apache.org> on 2023/06/15 19:22:50 UTC
[GitHub] [tvm] masahi opened a new pull request, #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
masahi opened a new pull request, #15111:
URL: https://github.com/apache/tvm/pull/15111
(no comment)
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1593602760
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* cc @billishyahao, @quic-sanirudh <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1614534204
I have compared against the exllama int4 kernel for decoding. It seems the FT GEMM kernel is overkill for mat x vec multiply - the custom mat x vec kernel from exllama is much faster.
```
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- --------- --------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
59.6 3,032,906,823 81,824 37,066.2 26,882.0 26,528 78,275 12,595.2 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
<cutlass::gemm::threadblock::DqMmaMultistage…
30.7 1,561,633,485 32,704 47,750.5 47,601.5 26,657 217,961 20,785.9 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
<cutlass::gemm::threadblock::DqMmaMultistage…
4.6 233,826,480 16,416 14,243.8 14,016.0 4,257 22,881 4,835.9 void attention_kernel_batched_impl<AttentionKernel<cutl
ass::half_t, cutlass::arch::Sm80, (bool)1, (…
1.2 59,620,306 513 116,218.9 116,196.0 114,597 120,773 660.2 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
<cutlass::gemm::threadblock::DqMmaMultistage…
```
```
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- --------- --------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
77.0 1,563,943,842 57,344 27,273.0 18,111.0 16,351 42,174 11,590.3 void q4_matmul_kernel<(bool)1, (bool)1, (bool)0>(const
__half *, const unsigned int *, __half *, co…
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1710616892
@LeiWang1999 yes, we have done similar perf comparison against cublas fp16, and indeed we found that the FT int4 is faster only for small M. This is not surprising since a larger M amortizes the cost of loading a large weight over each row (which may correspond to "batch" or "token" in llm inference) and int4 involves dequantize overhead.
The developer also talked about this point in their GTC23 presentation. Here is one of their slides.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] fxmarty commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "fxmarty (via GitHub)" <gi...@apache.org>.
fxmarty commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1614601193
Thank you, I'm trying to connect the dots between the different opinions I hear around using mma.sync.aligned op (which from what I understand of cutlass, the fastertransformer extension should come down to) vs not for GEMV, and your results are in line with my current intuition.
Now the only mystery remaining to me is why pytorch chooses a tensor core based kernel for GEMV in fp16 * fp16.
Maybe there could be opportunities to have a cutlass extension based rather on https://github.com/NVIDIA/cutlass/blob/main/include/cutlass/gemm/kernel/gemv.h / https://github.com/NVIDIA/cutlass/blob/main/include/cutlass/gemm/kernel/gemv_batched_strided.h / for the decoding, along with a gemm kernel for the prefill.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] sunggg commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "sunggg (via GitHub)" <gi...@apache.org>.
sunggg commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1710251218
@LeiWang1999 Thank you for sharing this. This is very interesting data point! May I ask how you collect these numbers? Did you exhaustively search over the possible tile combinations?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 merged pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 merged PR #15111:
URL: https://github.com/apache/tvm/pull/15111
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] cyang49 commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "cyang49 (via GitHub)" <gi...@apache.org>.
cyang49 commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1688780068
>
> I have compared against the exllama int4 kernel for decoding. It seems the FT GEMM kernel is overkill for mat x vec multiply - the custom mat x vec kernel from exllama is much faster.
>
> ```
> Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
>
> -------- --------------- --------- --------- --------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
> 59.6 3,032,906,823 81,824 37,066.2 26,882.0 26,528 78,275 12,595.2 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
> <cutlass::gemm::threadblock::DqMmaMultistage…
> 30.7 1,561,633,485 32,704 47,750.5 47,601.5 26,657 217,961 20,785.9 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
> <cutlass::gemm::threadblock::DqMmaMultistage…
> 4.6 233,826,480 16,416 14,243.8 14,016.0 4,257 22,881 4,835.9 void attention_kernel_batched_impl<AttentionKernel<cutl
> ass::half_t, cutlass::arch::Sm80, (bool)1, (…
> 1.2 59,620,306 513 116,218.9 116,196.0 114,597 120,773 660.2 void cutlass::Kernel<cutlass::gemm::kernel::GemmFpAIntB
> <cutlass::gemm::threadblock::DqMmaMultistage…
> ```
>
> ```
> Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
>
> -------- --------------- --------- --------- --------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
> 77.0 1,563,943,842 57,344 27,273.0 18,111.0 16,351 42,174 11,590.3 void q4_matmul_kernel<(bool)1, (bool)1, (bool)0>(const
> __half *, const unsigned int *, __half *, co…
> ```
@masahi what was the dimensions of the input and the GPU you tested this on? Thanks
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1233802790
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
Review Comment:
In general one might want to compile for a different GPU than the one in the compilation environment.
That said, if I can define a Python utility function with an optional sm argument, that can be called from TVMScript, that would be best. But is it possible to call an arbitrary python function from TVMScript?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1599411629
Any other comments? Otherwise, can someone merge this?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1709845769
> It seems that the performance of the default cutlass tile does not perform well on compute bound shape gemm
How do you know? The kernel uses a runtime heuristic to select a tile config. For now I'm hoping that this heuristic is "good enough".
https://github.com/tlc-pack/cutlass_fpA_intB_gemm/blob/main/cutlass_kernels/fpA_intB_gemm/fpA_intB_gemm_template.h#L818-L826
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] yzh119 commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "yzh119 (via GitHub)" <gi...@apache.org>.
yzh119 commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1233333463
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
Review Comment:
Can sm be inferred automatically? It doesn't look friendly to users.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232877450
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
+ sinfo_args=(R.Tensor((64, 64), dtype="int8"),),
+ )
+ lv3: R.Tensor((128,), dtype="float16") = lv[1]
+ lv6 = R.call_tir(
Review Comment:
I think it's a bit hard to define what "representative Q scheme" is going to be. Even for the simple one expected by FT, the weight needs to be packed in a very specific way.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1233809875
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
Review Comment:
Also CUTLASS users need to be very careful with the target sm anyway, so I hope it won't be a big deal in practice.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232880648
##########
src/runtime/contrib/cutlass/weight_preprocess.cc:
##########
@@ -0,0 +1,46 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <tvm/runtime/ndarray.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include "../../../3rdparty/cutlass_fpA_intB_gemm/cutlass_kernels/cutlass_preprocessors.h"
+
+namespace tvm {
+namespace runtime {
+
+TVM_REGISTER_GLOBAL("cutlass.ft_preprocess_weight_int4")
+ .set_body_typed([](NDArray packed_weight, int sm) {
+ int rows = packed_weight->shape[0];
+ int cols = packed_weight->shape[1];
+ std::vector<int8_t> input_cpu(rows * cols);
+ std::vector<int8_t> output_cpu(rows * cols);
+ packed_weight.CopyToBytes(input_cpu.data(), input_cpu.size());
+ // multiply cols by 2 since the "col" params in preprocess_weights refers to the column of
+ // the unpacked weight.
+ fastertransformer::preprocess_weights(output_cpu.data(), input_cpu.data(), rows, cols * 2,
Review Comment:
Note that this function is meant to be run at compile time. It needs to run on CPU since the preprocess functions are defined by FT in C++. I don't find them to be slow, so given the amount of effort required by GPU porting, I don't think it's worth it.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] sunggg commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "sunggg (via GitHub)" <gi...@apache.org>.
sunggg commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232777164
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
+ sinfo_args=(R.Tensor((64, 64), dtype="int8"),),
+ )
+ lv3: R.Tensor((128,), dtype="float16") = lv[1]
+ lv6 = R.call_tir(
Review Comment:
I'm wondering, in the future, once we settle down representative quantization schemes, would it make sense to introduce relax-level decode op?
##########
src/runtime/contrib/cutlass/weight_preprocess.cc:
##########
@@ -0,0 +1,46 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <tvm/runtime/ndarray.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include "../../../3rdparty/cutlass_fpA_intB_gemm/cutlass_kernels/cutlass_preprocessors.h"
+
+namespace tvm {
+namespace runtime {
+
+TVM_REGISTER_GLOBAL("cutlass.ft_preprocess_weight_int4")
+ .set_body_typed([](NDArray packed_weight, int sm) {
+ int rows = packed_weight->shape[0];
+ int cols = packed_weight->shape[1];
+ std::vector<int8_t> input_cpu(rows * cols);
+ std::vector<int8_t> output_cpu(rows * cols);
+ packed_weight.CopyToBytes(input_cpu.data(), input_cpu.size());
+ // multiply cols by 2 since the "col" params in preprocess_weights refers to the column of
+ // the unpacked weight.
+ fastertransformer::preprocess_weights(output_cpu.data(), input_cpu.data(), rows, cols * 2,
Review Comment:
Maybe good to add some documentation about why we need this preprocess, why it is happening on CPU, and if it is possible/better to move it on GPU?
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
+ sinfo_args=(R.Tensor((64, 64), dtype="int8"),),
+ )
+ lv3: R.Tensor((128,), dtype="float16") = lv[1]
+ lv6 = R.call_tir(
+ cls.decode, (lv2, lv3), out_sinfo=R.Tensor((64, 128), dtype="float16")
+ )
+ lv1_1: R.Tensor((64, 128), dtype="float16") = R.matmul(x, lv6, out_dtype="float16")
+ lv2_1: R.Tensor((64, 128), dtype="float16") = R.add(lv1_1, bias)
+ R.output(lv2_1)
+ return lv2_1
+
+ @R.function
+ def main_residual(
+ x: R.Tensor((64, 64), dtype="float16"),
+ residual: R.Tensor((64, 128), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 2})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
Review Comment:
If we always have to apply this preprocess, can we do the preprocess and then dump the weight?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232880648
##########
src/runtime/contrib/cutlass/weight_preprocess.cc:
##########
@@ -0,0 +1,46 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <tvm/runtime/ndarray.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include "../../../3rdparty/cutlass_fpA_intB_gemm/cutlass_kernels/cutlass_preprocessors.h"
+
+namespace tvm {
+namespace runtime {
+
+TVM_REGISTER_GLOBAL("cutlass.ft_preprocess_weight_int4")
+ .set_body_typed([](NDArray packed_weight, int sm) {
+ int rows = packed_weight->shape[0];
+ int cols = packed_weight->shape[1];
+ std::vector<int8_t> input_cpu(rows * cols);
+ std::vector<int8_t> output_cpu(rows * cols);
+ packed_weight.CopyToBytes(input_cpu.data(), input_cpu.size());
+ // multiply cols by 2 since the "col" params in preprocess_weights refers to the column of
+ // the unpacked weight.
+ fastertransformer::preprocess_weights(output_cpu.data(), input_cpu.data(), rows, cols * 2,
Review Comment:
It needs to run on CPU since the preprocess functions are defined by FT in C++. I don't find them to be slow, so given the amount of effort required by GPU porting, I don't think it's worth it.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1233802790
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
Review Comment:
In general one might want to compile for a different GPU than the one in the compilation environment. `cutlass.ft_preprocess_weight_int4` needs to be defined in `tvm_runtime` so we cannot do `Target::Current()` from there.
That said, if I can define a Python utility function with an optional sm argument, that can be called from TVMScript, that would be best. But is it possible to call an arbitrary python function from TVMScript?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] fxmarty commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "fxmarty (via GitHub)" <gi...@apache.org>.
fxmarty commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1606830801
Hi @masahi , I was wondering if you had benchmarked against e.g. available GPTQ cuda or triton-based kernels. Thank you!
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232886268
##########
src/runtime/contrib/cutlass/weight_preprocess.cc:
##########
@@ -0,0 +1,46 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+#include <tvm/runtime/ndarray.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include "../../../3rdparty/cutlass_fpA_intB_gemm/cutlass_kernels/cutlass_preprocessors.h"
+
+namespace tvm {
+namespace runtime {
+
+TVM_REGISTER_GLOBAL("cutlass.ft_preprocess_weight_int4")
+ .set_body_typed([](NDArray packed_weight, int sm) {
+ int rows = packed_weight->shape[0];
+ int cols = packed_weight->shape[1];
+ std::vector<int8_t> input_cpu(rows * cols);
+ std::vector<int8_t> output_cpu(rows * cols);
+ packed_weight.CopyToBytes(input_cpu.data(), input_cpu.size());
+ // multiply cols by 2 since the "col" params in preprocess_weights refers to the column of
+ // the unpacked weight.
+ fastertransformer::preprocess_weights(output_cpu.data(), input_cpu.data(), rows, cols * 2,
Review Comment:
Added some doc
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] LeiWang1999 commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "LeiWang1999 (via GitHub)" <gi...@apache.org>.
LeiWang1999 commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1710325858
hi, @sunggg code to reproduce:
- relax: https://gist.github.com/LeiWang1999/c3d644ba23a1c859c0aa2574f7412148
- cublas: https://gist.github.com/LeiWang1999/71779e5b4bc32fb63829420d42926b56
glad if it would help.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15111:
URL: https://github.com/apache/tvm/pull/15111#discussion_r1232880042
##########
tests/python/relax/test_codegen_cutlass.py:
##########
@@ -1250,5 +1250,243 @@ def main(
tvm.ir.assert_structural_equal(mod, Expected)
+def test_fp16A_int4B_gemm():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def decode(
+ A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ B: T.Buffer((T.int64(128),), "float16"),
+ decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ for i, j in T.grid(T.int64(64), T.int64(128)):
+ with T.block("decode"):
+ v_i, v_j = T.axis.remap("SS", [i, j])
+ T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
+ T.writes(decode_1[v_i, v_j])
+ decode_1[v_i, v_j] = (
+ T.Cast(
+ "float16",
+ T.shift_right(
+ T.shift_left(
+ T.bitwise_and(
+ T.shift_right(
+ T.Cast("int32", A[v_i, v_j // T.int64(2)]),
+ T.Cast("int32", v_j % T.int64(2)) * 4,
+ ),
+ 15,
+ ),
+ 28,
+ ),
+ 28,
+ ),
+ )
+ * B[v_j]
+ )
+
+ @T.prim_func
+ def encode(
+ A: T.Buffer((T.int64(128), T.int64(64)), "float16"),
+ w_gathered: T.Buffer((T.int64(64), T.int64(64)), "int8"),
+ compute: T.Buffer((T.int64(128),), "float16"),
+ ):
+ T.func_attr({"tir.noalias": T.bool(True)})
+ # with T.block("root"):
+ max_abs_value = T.alloc_buffer((T.int64(128),), "float16")
+ scale = T.alloc_buffer((T.int64(128),))
+ for i, k in T.grid(T.int64(128), T.int64(64)):
+ with T.block("max_abs_value"):
+ v_i, v_k = T.axis.remap("SR", [i, k])
+ T.reads(A[v_i, v_k])
+ T.writes(max_abs_value[v_i])
+ with T.init():
+ max_abs_value[v_i] = T.float16(-65504)
+ max_abs_value[v_i] = T.max(max_abs_value[v_i], T.fabs(A[v_i, v_k]))
+ for i in range(T.int64(128)):
+ with T.block("scale"):
+ v_i = T.axis.spatial(T.int64(128), i)
+ T.reads(max_abs_value[v_i])
+ T.writes(scale[v_i])
+ scale[v_i] = T.max(
+ T.Cast("float32", max_abs_value[v_i]), T.float32(0.0001)
+ ) * T.float32(0.125)
+ for j, i, k in T.grid(T.int64(64), T.int64(64), T.int64(2)):
+ with T.block("w_gathered"):
+ v_j, v_i, v_k = T.axis.remap("SSR", [j, i, k])
+ T.reads(A[v_i * T.int64(2) + v_k, v_j], scale[v_i * T.int64(2) + v_k])
+ T.writes(w_gathered[v_j, v_i])
+ with T.init():
+ w_gathered[v_j, v_i] = T.int8(0)
+ w_gathered[v_j, v_i] = T.bitwise_or(
+ w_gathered[v_j, v_i],
+ T.if_then_else(
+ v_i * T.int64(2) + v_k < T.int64(128),
+ T.shift_left(
+ T.bitwise_and(
+ T.Cast(
+ "int8",
+ T.min(
+ T.max(
+ T.round(
+ T.Cast(
+ "float32", A[v_i * T.int64(2) + v_k, v_j]
+ )
+ / scale[v_i * T.int64(2) + v_k]
+ ),
+ T.float32(-8),
+ ),
+ T.float32(7),
+ ),
+ ),
+ T.int8(15),
+ ),
+ T.Cast("int8", v_k) * T.int8(4),
+ ),
+ T.int8(0),
+ ),
+ )
+ for i0 in range(T.int64(128)):
+ with T.block("compute"):
+ v_i0 = T.axis.spatial(T.int64(128), i0)
+ T.reads(scale[v_i0])
+ T.writes(compute[v_i0])
+ compute[v_i0] = T.Cast("float16", scale[v_i0])
+
+ @R.function
+ def main_bias(
+ x: R.Tensor((64, 64), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 1})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
+ lv1,
+ 80,
+ sinfo_args=(R.Tensor((64, 64), dtype="int8"),),
+ )
+ lv3: R.Tensor((128,), dtype="float16") = lv[1]
+ lv6 = R.call_tir(
+ cls.decode, (lv2, lv3), out_sinfo=R.Tensor((64, 128), dtype="float16")
+ )
+ lv1_1: R.Tensor((64, 128), dtype="float16") = R.matmul(x, lv6, out_dtype="float16")
+ lv2_1: R.Tensor((64, 128), dtype="float16") = R.add(lv1_1, bias)
+ R.output(lv2_1)
+ return lv2_1
+
+ @R.function
+ def main_residual(
+ x: R.Tensor((64, 64), dtype="float16"),
+ residual: R.Tensor((64, 128), dtype="float16"),
+ y: R.Tensor((128, 64), dtype="float16"),
+ bias: R.Tensor((1, 128), dtype="float16"),
+ ) -> R.Tensor((64, 128), dtype="float16"):
+ R.func_attr({"num_input": 2})
+ cls = Module
+ with R.dataflow():
+ lv = R.call_tir(
+ cls.encode,
+ (y,),
+ out_sinfo=[R.Tensor((64, 64), dtype="int8"), R.Tensor((128,), dtype="float16")],
+ )
+ lv1 = lv[0]
+ lv2 = R.call_pure_packed(
+ "cutlass.ft_preprocess_weight_int4",
Review Comment:
Yes that's possible, encode + preprocess and dump. In practice, if we do `LiftTransformParams` pass, this preprocess is already done as part of running `transform_params` function (together with encode). So we are already doing what you described.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1688813871
This was the profiler output on one e2e inference on vicuna 7B decoder. The GPU was RTX 4080.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] LeiWang1999 commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "LeiWang1999 (via GitHub)" <gi...@apache.org>.
LeiWang1999 commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1709871632
thanks, I see. I mentioned that because I benchmarked the performance on the Llama-70b GEMM (cuBLAS FP16xFP16 vs. Relax.Cutlass fp16xint4).
<html xmlns:v="urn:schemas-microsoft-com:vml"
xmlns:o="urn:schemas-microsoft-com:office:office"
xmlns:x="urn:schemas-microsoft-com:office:excel"
xmlns="http://www.w3.org/TR/REC-html40">
<body link="#467886" vlink="#96607D">
M | N | K | cublas | cutlass-fpa-intb | speedup
-- | -- | -- | -- | -- | --
1 | 1024 | 8192 | 0.030588 | 0.046896935 | 0.652236035
1 | 8192 | 8192 | 0.199339 | 0.087285042 | 2.28376673
1 | 8192 | 28672 | 1.055949 | 0.187683105 | 5.626232851
1 | 28672 | 8192 | 0.672551 | 0.192594528 | 3.492055583
16 | 1024 | 8192 | 0.036352 | 0.095558167 | 0.380417524
16 | 8192 | 8192 | 0.19456 | 0.086522102 | 2.248674049
16 | 8192 | 28672 | 0.666054 | 0.187945366 | 3.54386972
16 | 28672 | 8192 | 0.669696 | 0.19159317 | 3.495406297
32 | 1024 | 8192 | 0.037856 | 0.101447105 | 0.373159996
32 | 8192 | 8192 | 0.196592 | 0.101852417 | 1.930165321
32 | 8192 | 28672 | 0.664686 | 0.319719315 | 2.078966443
32 | 28672 | 8192 | 0.67072 | 0.279974937 | 2.395642936
64 | 1024 | 8192 | 0.05376 | 0.170302391 | 0.31567378
64 | 8192 | 8192 | 0.199168 | 0.175619125 | 1.134090585
64 | 8192 | 28672 | 0.675594 | 0.57182312 | 1.181474212
64 | 28672 | 8192 | 0.681984 | 0.283193588 | 2.408190141
128 | 1024 | 8192 | 0.078336 | 0.051283836 | 1.527498838
128 | 8192 | 8192 | 0.238592 | 0.178599358 | 1.335906254
128 | 8192 | 28672 | 0.739888 | 0.575089455 | 1.286561606
128 | 28672 | 8192 | 0.714752 | 0.521111488 | 1.371591367
1024 | 1024 | 8192 | 0.28323 | 0.171804428 | 1.648558319
1024 | 8192 | 8192 | 1.158315 | 1.086783409 | 1.065819275
1024 | 8192 | 28672 | 4.166997 | 4.37412262 | 0.952647604
1024 | 28672 | 8192 | 3.709171 | 3.643035889 | 1.01815373
4096 | 1024 | 8192 | 0.597701 | 0.690698624 | 0.865357094
4096 | 8192 | 8192 | 4.334677 | 6.313800812 | 0.686540065
4096 | 8192 | 28672 | 16.15016 | 23.39763641 | 0.690247328
4096 | 28672 | 8192 | 14.95114 | 22.15902805 | 0.674720217
8192 | 1024 | 8192 | 1.200128 | 1.547074318 | 0.775740341
8192 | 8192 | 8192 | 8.550848 | 12.74256706 | 0.671045949
8192 | 8192 | 28672 | 30.88835 | 44.36366558 | 0.696253236
8192 | 28672 | 8192 | 30.19989 | 44.58220005 | 0.677397932
16384 | 1024 | 8192 | 2.349141 | 3.225851059 | 0.728223751
16384 | 8192 | 8192 | 16.97864 | 25.81973076 | 0.657583766
16384 | 8192 | 28672 | 61.27206 | 91.48414135 | 0.669756127
16384 | 28672 | 8192 | 61.03514 | 90.86410999 | 0.671718869
</body>
</html>
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] LeiWang1999 commented on pull request #15111: [Unity][BYOC] Integrate fp16 A - int4 B GEMM kernel from FasterTransformer into CUTLASS BYOC
Posted by "LeiWang1999 (via GitHub)" <gi...@apache.org>.
LeiWang1999 commented on PR #15111:
URL: https://github.com/apache/tvm/pull/15111#issuecomment-1709786106
It seems that the performance of the default cutlass tile does not perform well on compute bound shape gemm. Do we have any plans to support tile tuning for this backend?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org