Skip to content

Official support for mgpu vectorization hints#405

Merged
neoblizz merged 2 commits intomainfrom
neoblizz/vec-hints
Feb 28, 2026
Merged

Official support for mgpu vectorization hints#405
neoblizz merged 2 commits intomainfrom
neoblizz/vec-hints

Conversation

@neoblizz
Copy link
Member

Motivation

This pull request adds support for vectorization hints to various distributed memory operations in the iris.py module. By introducing an optional hint parameter, the code now allows callers to specify alignment and vectorization preferences for pointer translation and memory access, which can lead to more efficient memory operations on supported hardware.

The most important changes are:

Vectorization hint support:

  • Added an optional hint parameter (of type tl.constexpr) to the __translate function and all distributed memory operation methods (e.g., load, store, copy, get, and all atomic operations) in iris.py. This parameter allows callers to specify alignment and vectorization preferences for pointer translation. [1] [2] [3] [4] [5] [6] [7] [8] [9] [10] [11] [12] [13] [14] [15] [16]
  • Updated the pointer translation logic in __translate, copy, and related functions to apply the vectorization hint using tl.multiple_of and tl.max_contiguous when the hint is provided. [1] [2] [3]

Documentation updates:

  • Updated docstrings for all affected functions to describe the new hint parameter, including usage examples and explanations of how to specify the hint for 1-D and N-D cases. [1] [2] [3] [4] [5]

These changes provide more control over memory access patterns, which can help optimize performance for distributed and vectorized workloads.

Submission Checklist

Copilot AI review requested due to automatic review settings February 28, 2026 13:30
@github-actions github-actions bot added in-progress We are working on it iris Iris project issue labels Feb 28, 2026
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Adds official mgpu vectorization hint support across Iris distributed-memory helpers, enabling callers to provide alignment/contiguity assumptions to improve generated memory code where valid.

Changes:

  • Added an optional hint: tl.constexpr parameter to pointer translation and all distributed memory ops (load/store/copy/get/put/atomics).
  • Applied tl.multiple_of + tl.max_contiguous to translated pointers when hint is provided.
  • Expanded docstrings to document the new hint parameter on the public APIs.

Copy link
Collaborator

@mawad-amd mawad-amd left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let's introduce one new argument per hint and add tests.

@neoblizz
Copy link
Member Author

Let's introduce one new argument per hint and add tests.

See the comment above.

The current local tests I was doing was assembly generation and inspection.

@neoblizz neoblizz merged commit a3fca93 into main Feb 28, 2026
77 of 78 checks passed
@neoblizz neoblizz deleted the neoblizz/vec-hints branch February 28, 2026 19:05
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

in-progress We are working on it iris Iris project issue

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants