cbalint13 commented on code in PR #13621:
URL: https://github.com/apache/tvm/pull/13621#discussion_r1056822933
##########
python/tvm/tir/tensor_intrin/x86.py:
##########
@@ -67,8 +68,50 @@ def dot_product_16x4_u8i8i32_vnni(
)
+mem_scope="global"
[email protected]_func
+def dot_product_16x4_u8i8i32_avx512(a: T.handle, b: T.handle, c: T.handle,) ->
None:
+ A = T.match_buffer(a, (4,), "uint8", offset_factor=1, scope=mem_scope)
+ B = T.match_buffer(b, (16, 4), "int8", offset_factor=1, scope=mem_scope)
+ C = T.match_buffer(c, (16,), "int32", offset_factor=1, scope=mem_scope)
+ with T.block("root"):
+ T.reads(C[0:16], A[0:4], B[0:16, 0:4])
+ T.writes(C[0:16])
+
+ A_u8x4 = A.vload([0], "uint8x4")
+ A_i32 = T.reinterpret(A_u8x4, dtype="int32")
+ A_brdcst = T.broadcast(A_i32, 16)
+ A_u8x64 = T.reinterpret(A_brdcst, dtype="uint8x64")
+
+ B_i8x64 = B.vload([0, 0], dtype="int8x64")
+
+ Red = T.call_llvm_pure_intrin(
+ T.llvm_lookup_intrinsic_id("llvm.x86.avx512.pmaddubs.w.512"),
+ T.uint32(2),
+ A_u8x64,
+ B_i8x64,
+ dtype="int16x32",
+ )
+
+ One_int16x32 = tvm.tir.const(1, "int16x32")
Review Comment:
* ```One_int16x32 = T.int16x32(1)``` ?
* Would remove ```import tvm``` line.
##########
python/tvm/tir/tensor_intrin/x86.py:
##########
@@ -67,8 +68,50 @@ def dot_product_16x4_u8i8i32_vnni(
)
+mem_scope="global"
[email protected]_func
+def dot_product_16x4_u8i8i32_avx512(a: T.handle, b: T.handle, c: T.handle,) ->
None:
+ A = T.match_buffer(a, (4,), "uint8", offset_factor=1, scope=mem_scope)
+ B = T.match_buffer(b, (16, 4), "int8", offset_factor=1, scope=mem_scope)
+ C = T.match_buffer(c, (16,), "int32", offset_factor=1, scope=mem_scope)
Review Comment:
[cosmetic] Can be similar to the ```vnni``` part, as below ?
https://github.com/apache/tvm/blob/209845fb910aac76f1388d868bb4fe4a46d9170f/python/tvm/tir/tensor_intrin/x86.py#L44-L49
Also, scope ```global``` is default anyway.
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
