Skip to content

Lecture 14 weird result with TRITON_INTERPRET=1 #51

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
noklam opened this issue Apr 16, 2025 · 1 comment · May be fixed by #52
Open

Lecture 14 weird result with TRITON_INTERPRET=1 #51

noklam opened this issue Apr 16, 2025 · 1 comment · May be fixed by #52

Comments

@noklam
Copy link

noklam commented Apr 16, 2025

Image

The matmul result is obviously wrong. Then I execute the next cell with (512,512), which produce a correct result surprisingly.

Image

I suspect there are some bug with the index or missing guard clause. See another example with (16,16), which match the batch size also passed.

Image

p.s. I do not have a GPU and running this on a Mac so everything is run in CPU mode and TRITON_INTERPRET=1

@noklam
Copy link
Author

noklam commented Apr 16, 2025

Reading the code again, should mask be applied in tl.load so that the dot product only calculate element within the bound?

Adding the mask will pass all 3 cases.

    # Get 1d mask for 
    for _ in range(0, k, bk):
        mask_a = get_2d_mask(rm, rn, m, k) # add mask
        mask_b = get_2d_mask(rm, rn, k, n)  # add mask
            
        a = tl.load(offs_a, mask_a)
        b = tl.load(offs_b, mask_b)
        acc += tl.dot(a, b, allow_tf32=False) # matmul in block ; Weirdness: allow_tf32 must be set to False for older GPUs, otherwise won't compile
        
        # print_if(f"pid: {pid_m, pid_n} |\na: {a} | \nb: {b}", "")
        
        # increase offets, so next iteration loads next chunks
        offs_a += bk * stride_ak
        offs_b += bk * stride_bk

@noklam noklam linked a pull request Apr 16, 2025 that will close this issue
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging a pull request may close this issue.

1 participant