Skip to content

[PTX] Enable migration of ldmatrix #2692

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
May 7, 2025

Conversation

TejaX-Alaghari
Copy link
Contributor

This PR adds support for the migration of ldmatrix PTX ASM API

@TejaX-Alaghari TejaX-Alaghari requested a review from a team as a code owner March 6, 2025 17:29
@zhiweij1 zhiweij1 changed the title [PRX] Enable migration of ldmatrix [PTX] Enable migration of ldmatrix Mar 7, 2025
@TejaX-Alaghari TejaX-Alaghari force-pushed the ptx branch 5 times, most recently from bd88679 to 4492ec9 Compare March 13, 2025 05:16
Copy link
Contributor

@tomflinda tomflinda left a comment

Choose a reason for hiding this comment

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

LGTM

@TejaX-Alaghari TejaX-Alaghari force-pushed the ptx branch 2 times, most recently from de49c1d to 7f533ae Compare March 23, 2025 05:41
@@ -2055,6 +2055,103 @@ class joint_matrix {
matrix_accessor x;
const size_t num_elements;
};

/// Loads 1 8x8 b16 matrix from shared memory to local memory (32-bits per wi)
Copy link
Contributor

Choose a reason for hiding this comment

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

from local memory to private memory of each work item?

}
}

/// Loads 2 8x8 b16 matrix from shared memory to local memory (32-bits per wi)
Copy link
Contributor

Choose a reason for hiding this comment

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

Please use sycl language in the comments.

@@ -2055,6 +2055,98 @@ class joint_matrix {
matrix_accessor x;
const size_t num_elements;
};

/// Loads 1 8x8 b16 matrix from local memory to private memory (32-bits per wi)
Copy link
Contributor

Choose a reason for hiding this comment

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

Suggested change
/// Loads 1 8x8 b16 matrix from local memory to private memory (32-bits per wi)
/// Loads 1 8x8 b16 matrix from local memory to private memory (32-bits per work item)

Copy link
Contributor

Choose a reason for hiding this comment

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

Please give more details on how the loaded data will be distributed on private memory. is it a work-group function or sub-group function? Make sure user can use it after reading the comments.

@TejaX-Alaghari TejaX-Alaghari force-pushed the ptx branch 2 times, most recently from 0463353 to 63ffbaf Compare April 29, 2025 10:08
@TejaX-Alaghari TejaX-Alaghari force-pushed the ptx branch 4 times, most recently from 22018c1 to df279d5 Compare April 30, 2025 09:28
@zhimingwang36 zhimingwang36 merged commit 971e72b into oneapi-src:SYCLomatic May 7, 2025
7 checks passed
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.

5 participants