Feature Description
Unify the variable names for the alloc buffer.
Proposed Solution
The alloc_ub flag is set to shared.ub, the alloc_l1 flag is set to shared.l1, and the alloc_shared flag is set to shared. Then, the pass flag is changed to shared.ub or shared.l1.
|
def alloc_shared(shape, dtype, scope="shared.dyn"): |
|
"""Allocate a shared memory buffer for inter-thread communication. |
|
|
|
Args: |
|
shape (tuple): The shape of the buffer to allocate |
|
dtype (str): The data type of the buffer (e.g., 'float32', 'int32') |
|
scope (str, optional): The memory scope. Defaults to "shared.dyn" |
|
|
|
Returns: |
|
T.Buffer: A TVM buffer object allocated in shared memory |
|
""" |
|
if dtype == "bool": |
|
# lei: This is a hack to handle bool type. |
|
# Because tilelang's merge smem pass cannot merge bool type currently. |
|
scope = "shared" |
|
return T.alloc_buffer(shape, dtype, scope=scope) |
|
def alloc_L1(shape, dtype): |
|
return T.alloc_buffer(shape, dtype, scope="shared.l1") |
|
|
|
|
|
def alloc_L0A(shape, dtype): |
|
return T.alloc_buffer(shape, dtype, scope="wmma.matrix_a") |
|
|
|
|
|
def alloc_L0B(shape, dtype): |
|
return T.alloc_buffer(shape, dtype, scope="wmma.matrix_b") |
|
|
|
|
|
def alloc_L0C(shape, dtype): |
|
return T.alloc_buffer(shape, dtype, scope="wmma.accumulator") |
|
|
|
|
|
def alloc_ub(shape, dtype): |
|
return T.alloc_buffer(shape, dtype, scope="shared") |
Feature Type
Feature Description
Unify the variable names for the alloc buffer.
Proposed Solution
The alloc_ub flag is set to shared.ub, the alloc_l1 flag is set to shared.l1, and the alloc_shared flag is set to shared. Then, the pass flag is changed to shared.ub or shared.l1.
tilelang-ascend/tilelang/language/allocate.py
Lines 32 to 47 in 35732a5
tilelang-ascend/tilelang/language/allocate.py
Lines 154 to 171 in 35732a5
Feature Type