NVIDIA / NVIDIA/cutlass

[QST]Question about LdMatrix16x16x8bOp

Open
#2,871 3 comments 0 reactions 0 assignees View on GitHub

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.

Image

https://github.com/NVIDIA/cutlass/blob/49bd6bf1ba80abd588a56a3f3af2f1bcd41d215f/python/CuTeDSL/cutlass/cute/nvgpu/warp/copy.py#L82C1-L114C73

@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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.