cbalint13 commented on code in PR #13642:
URL: https://github.com/apache/tvm/pull/13642#discussion_r1058846644
##########
python/tvm/topi/x86/tensor_intrin.py:
##########
@@ -348,3 +348,227 @@ def _instr(index):
binds={data: a_buffer, kernel: b_buffer},
default_buffer_params=buffer_params,
)
+
+
+def dot_32x128x32_u8s8s32_sapphirerapids(LDA):
+ """
+ Int8 dot product by every 16x64 elements using AMX-TMUL Sapphire Rapids
instructions.
+ The tdpxxd instruction takes two tile of uint8 and int8 datatype --
data[16][64] and
+ kernel[1][16][16][4] -- and computes a dot product of data[16][16] in
int32 datatype.
+
+ (Physically, to efficiently leveraging the tile register, we constructing
a 2x2 tiles
+ matmul which performs 32x128x32 in total)
+
+ The pseudo code is as follows:
+ for(k=0; k<2; k++){
+ for(n=0; n<2; n++){
+ tileload64(tmm_b, B)
+ for(m=0; m<2; m++){
+ if(n==0)
+ tileload64(tmm_a, A)
Review Comment:
@masahi ,
@Qianshui-Jiang ,
Don't want hold up the PR, I think it is fine to merge with this startup AMX
implementation.
I leave here a possible way to workaround the rigidity of LLVM's AMX
related intrinsic declaration:
* Use ```tvm.tir.call_extern``` insead of ```tvm.tir.call_llvm_intrin```
like
[here](https://github.com/apache/tvm/blob/8551a5c71faaa2eedc8b1c49a39be0fe4688c021/python/tvm/topi/arm_cpu/mprofile/dsp/micro_kernel/gemm.py#L107-L109)
.
* You can than pass argument (e.g. register number, etc ...) into the custom
called function.
* You can use the argument inside that function and emit the AMX ISA as
bytes: [ __asm(.word
packed())](https://github.com/cbalint13/OLIMP/blob/main/demo/icebreaker/src/rv32-custom.h#L52-L53)
This is how it is done to emit custom part of RISC-V ISA, that is also too
rigid using just the LLVM declarations.
--
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]
