tilelang.intrinsics.mma_sp_layout¶

Functions¶

Module Contents¶

tilelang.intrinsics.mma_sp_layout.shared_16x16_to_mma_sp_layout_sr_a(i, j)¶
tilelang.intrinsics.mma_sp_layout.shared_16x16_to_mma_sp_layout_sr_b(i, j)¶
tilelang.intrinsics.mma_sp_layout.shared_16x32_to_mma_sp_layout_sr_a(i, j)¶
tilelang.intrinsics.mma_sp_layout.shared_16x32_to_mma_sp_layout_sr_b(i, j)¶
tilelang.intrinsics.mma_sp_layout.shared_16x64_to_mma_sp_layout_sr_a(i, j)¶
tilelang.intrinsics.mma_sp_layout.shared_16x64_to_mma_sp_layout_sr_b(i, j)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_a_32x4_to_shared_16x16_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_a_32x8_to_shared_16x32_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_a_32x16_to_shared_16x64_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_b_32x8_to_shared_16x16_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_b_32x16_to_shared_16x32_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.mma_sp_load_b_32x32_to_shared_16x64_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.get_logical_id_32bit(thread_id)¶
Parameters:

thread_id (int)

Return type:

int

tilelang.intrinsics.mma_sp_layout.metadata_8bit_load_32x4_to_shared_16x4_layout_32bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.metadata_16bit_load_32x2_to_shared_16x2_layout_32bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.metadata_8bit_load_32x4_to_shared_16x4_layout_16bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.metadata_16bit_load_32x2_to_shared_16x2_layout_16bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.get_logical_id_8bit(thread_id)¶
Parameters:

thread_id (int)

Return type:

int

tilelang.intrinsics.mma_sp_layout.metadata_8bit_load_32x4_to_shared_16x4_layout_8bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.metadata_16bit_load_32x2_to_shared_16x4_layout_8bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.metadata_32bit_load_32x1_to_shared_16x2_layout_8bit(thread_id, local_id)¶
Parameters:
  • thread_id (int)

  • local_id (int)

Return type:

tuple[int, int]

tilelang.intrinsics.mma_sp_layout.ldmatrix_trans_32x8_to_shared_16x16_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.ldmatrix_32x16_to_shared_32x16_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.ldmatrix_trans_32x16_to_shared_16x32_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.ldmatrix_trans_32x32_to_shared_shared_16x64_layout(thread_id, local_id)¶
tilelang.intrinsics.mma_sp_layout.get_ldmatrix_offset_b(matrix, row_idx, col_idx, stride, dtype='float16', transposed=False)¶
Parameters:
  • matrix (Literal['B'])

  • dtype (Literal['float16', 'int8'])

  • transposed (bool)