Skip to content

Conversation

@Jokeren
Copy link
Contributor

@Jokeren Jokeren commented Oct 23, 2025

No description provided.

return self.builder.get_shared_bank_conflicts(reg_attr, shared_attr, list(distr_ty.shape),
distr_ty.element_ty.primitive_bitwidth)

def to_linear(self, layout, shape):
Copy link
Contributor Author

Choose a reason for hiding this comment

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

To be rigorous, we call them "encodings" in TritonGPU instead of "layouts", but somehow gluon uses the term "layout". So not sure how to differentiate between "encodings" and the underlying more generic "layouts" in terminology

Copy link
Contributor

Choose a reason for hiding this comment

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

In my head, an encoding is a map from hardware into a tensor. A linear layouts is more generic, you can have an LL that represents a conversion from hardware to hardware calling invertAndCompose and so on

Copy link
Contributor Author

Choose a reason for hiding this comment

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

an encoding is a map from hardware into a tensor

Yes, so I think technically what we defined in gluon are "encodings", right?

Copy link
Contributor

Choose a reason for hiding this comment

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

yep! But again, not the biggest issue, it's more of a technicality I think

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Sure, we have two methods here, (1) convert a gluon encoding to linear encoding, (2) convert gluon encoding to linear layout. The former I named "to_linear", while the latter I named "to_linear_layout". I don't think they are good names so wanted to ask you about.

from . import language
from . import testing
from . import tools
from triton._C.libtriton.linear_layout import LinearLayout
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Is it the right way to import?

@@ -0,0 +1,77 @@
from triton import LinearLayout
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Should I just rewrite all cases from LinearLayoutTests in Python?

Copy link

@chatgpt-codex-connector chatgpt-codex-connector bot left a comment

Choose a reason for hiding this comment

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

💡 Codex Review

Here are some automated review suggestions for this pull request.

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment on lines 202 to 210
.def("get_matrix_view", [](const LinearLayout &self) {
std::unique_ptr<uint64_t[]> matrix = mlir::triton::getMatrix(self);
auto nRows = self.getNumOutDims();
auto nCols = self.getNumInDims();
std::vector<std::vector<int>> result(nRows, std::vector<int>(nCols));
for (size_t i = 0; i < nRows; ++i) {
for (size_t j = 0; j < nCols; ++j) {
result[i][j] = (matrix[i] >> j) & 1;
}

Choose a reason for hiding this comment

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

P1 Badge Use total bit widths when exporting matrix view

The new Python binding for LinearLayout.get_matrix_view iterates self.getNumOutDims() rows and self.getNumInDims() columns, but mlir::triton::getMatrix allocates its matrix with getTotalOutDimSizeLog2() rows and getTotalInDimSizeLog2() columns. When an output or input dimension has size 1, getMatrix returns a zero‑sized array while the wrapper still dereferences matrix[0], which is undefined behaviour and can crash. Even when the layout is valid, returning only the number of dimensions drops rows/columns for multi‑bit dimensions, so the exposed matrix is incorrect. Use the total bit counts to size the loops and result array before reading from matrix.

Useful? React with 👍 / 👎.

return result;
},
py::arg("inputs"))
.def("get_matrix_view", [](const LinearLayout &self) {
Copy link
Contributor Author

Choose a reason for hiding this comment

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

convert to numpy to allow validating matrix operations

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants