triton-lang / triton

Development repository for the Triton language and compiler
https://triton-lang.org/
MIT License
13.57k stars 1.67k forks source link

about llvm.struct and vector #5091

Open Shaquille-Wu opened 3 weeks ago

Shaquille-Wu commented 3 weeks ago

Hello, triton experts. I have a question about llvm.struct and vector

I think the element type of llvm.struct cannot be ensured the same, but, the each element type of vector must be the same with each other.

So, I guess vector has more advantages than llvm.struct, if llvm execute optimization pass(only guess)

In triton, I found output of loadOp is llvm.struct if the loadOp is vectorized. and the code like this:

%36 = llvm.inline_asm has_side_effects asm_dialect = att operand_attrs = [] "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b" %34, %32 : (!llvm.ptr<1>, i1) -> !llvm.struct<(i32, i32, i32, i32)> loc(#loc8)

I don't know what is the reason, it is llvm.struct instead of vector. why?

Is there anyone would like to teach me?

minjang commented 3 weeks ago

Your point is valid. I tried to change llvm.struct into vector for numerical types and LLVMFixedVectorType for pointer types. It works okay, but you'd need to change several places. While llvm.struct has a problem you pointed, it's handy as it can take any types. To address your concern, we could simply add an assert before/after unpack/packLLElements.