Testing triton when C matrix has some initial values. Possible?

[navdeepkk]

Hi, Is it possible to benchmark triton when C matrix is not assumed to be all zeros, i.e., it is loaded from global memory rather than initializing an accumulator tile with zeros?

Thanks!

[navdeepkk]

@ptillet any help would be appreciated. Thanks!

[ptillet]

@navdeepkk Sorry for the delay. I was double-checking something. So right now it seems like initialize matmul accumulator with loaded values is buggy. But what you can do instead is to add an accumulation after the matmul loop.

acc = 0
for k in range(K, 0, -BLOCK_K):
  ...
D = ... # construct pointer to the other tensor
acc += tl.load(D)

I will work on fixing bugs so that one can do

D = ... # construct pointer to the other tensor
acc = tl.load(D)
for k in range(K, 0, -BLOCK_K):
  ...

instead

[navdeepkk]

Thanks! I'll try this.

[ptillet]

Actually I just checked and this doesn't seem to work anymore now ... But I'll be working on it over the next few days.

[navdeepkk]

Oh Okay, I have some matmul kernels generated, which load C from global memory. I wanted to compare against Triton. I'll get on that when this is fixed. Thanks!

[ptillet]

Hey. This should work now on top of master :) Just add for example c = c + tl.load(c_ptrs, mask=c_mask) there: https://github.com/openai/triton/blob/master/python/tutorials/03-matrix-multiplication.py#L253. You can also initialize the accumulator here https://github.com/openai/triton/blob/master/python/tutorials/03-matrix-multiplication.py#L228, but it should be less efficient, and for now it'll only work if you explicitly cast it to float32 after loading. Let me know if you have any more issue