Skip to content
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

[TL] Enhance TL to import customized c headers #179

Merged
merged 48 commits into from
Sep 7, 2024

Conversation

LeiWang1999
Copy link
Contributor

Introduce two primitives. The first example of injection:

header = """
__device__ void acquire_lock(int* lock) {
    while (atomicCAS(lock , 1, 1) != 1) {
        // Busy-wait (spin-wait) until the lock is acquired
    }
}
"""

with T.Kernel(total_sm, threads=threads, prelude=header) as pid:
        T.call_extern(
                        "handle",
                        "atomicExch",
                        T.address_of(locks_buf[tile_id]),
                        1,
                    )

Or with T.import_source

with T.Kernel(total_sm, threads=threads) as pid:

            A_shared = T.alloc_shared(A_shared_shape, dtypeAB, scope=shared_scope)
            B_shared = T.alloc_shared(B_shared_shape, dtypeAB, scope=shared_scope)
            C_shared = T.alloc_shared(C_shared_shape, dtypeC, scope=shared_scope)
            A_local = T.alloc_fragment((warp_rows * local_size), dtypeAB, scope="local")
            B_local = T.alloc_fragment((warp_cols * local_size), dtypeAB, scope="local")
            C_local = T.alloc_fragment((warp_rows * warp_cols * local_size), accum_dtype, scope="local")
            start_iter = T.alloc_fragment((1,), "int32", "local")
            end_iter = T.alloc_fragment((1,), "int32", "local")
            T.import_source(header)

            T.annotate_layout(
                {
                    A_shared: make_swizzle_layout(A_shared),
                    B_shared: make_swizzle_layout(B_shared),
                }
            )

@LeiWang1999 LeiWang1999 merged commit 2f6d316 into microsoft:main Sep 7, 2024
6 checks passed
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 this pull request may close these issues.

1 participant