[QST]Question about LdMatrix16x16x8bOp
Open
Nobody has claimed this yet.
? - Needs Triage
inactive-30d
inactive-90d
question
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Hi, I have a question about the class LdMatrix16x16x8bOp. Transpose should not be asserted to be true, since PTX supports {.trans} modifier for ldmatrix.
@dataclass(frozen=True)
class LdMatrix16x16x8bOp(BaseOp):
"""
16x16 8-bit ``ldmatrix`` Operation.
See the `PTX documentation <https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-load-instruction-ldmatrix>`__.
This operation corresponds to the ``.m16n16`` and the ``.b16`` qualifiers.
"""
def __init__(self, num_matrices: int) -> None:
super().__init__(transpose=True, num_matrices=num_matrices)
self._verify()
def _verify(self):
assert self.transpose, "transpose must be True"
if self.num_matrices not in [1, 2]:
raise OpError(
self,
"expects the 'num_matrices' Op parameter to be one of [1,2]",
)
def _make_trait(
self, copy_internal_type: Type[Numeric], *, loc=None, ip=None, **kwargs
) -> "LdMatrix16x16x8bTrait":
mode = _pack_shape((16, 16), loc=loc, ip=ip)
ty = _cute_nvgpu_ir.CopyAtomLdsmType.get(
copy_internal_type.mlir_type,
mode.type.attribute,
_cute_nvgpu_ir.LdsmSzPattern.u8,
self.num_matrices,
ir.UnitAttr.get(),
)
return LdMatrix16x16x8bTrait(cute.make_atom(ty, loc=loc, ip=ip))
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start in python/CuTeDSL/cutlass/cute/nvgpu/warp/copy.py at LdMatrix16x16x8bOp, especially its constructor and _verify method. Read the linked PTX ldmatrix documentation to confirm the supported transpose modifiers, then determine the expected behavior for both forms; the issue is done when the operation no longer rejects a PTX-supported modifier and its validation remains correct.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- hpc
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 38/100