Skip to content

MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform#1414

Merged
yingluAMD merged 42 commits into
ROCm:developfrom
yingluAMD:conv_tf32_poc
Nov 16, 2025
Merged

MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform#1414
yingluAMD merged 42 commits into
ROCm:developfrom
yingluAMD:conv_tf32_poc

Conversation

@yingluAMD
Copy link
Copy Markdown
Contributor

@yingluAMD yingluAMD commented Sep 1, 2025

Motivation

CDNA3 series support TF32 in matrix core natively. While TF32 is not supported in MIOpen now. This PR is a POC of enabling TF32.

Technical Details

All kernel is changed in CK( PR:2763). Below are the changes in miopen:

  • Change problem and kernel instance invoker to invoker TF32 kernel in CK.
  • Add environment to control whether use TF32.

Test Plan

Test Result

Submission Checklist

@yingluAMD yingluAMD requested a review from a team as a code owner September 1, 2025 07:24
@yingluAMD yingluAMD changed the title feature:tf32:demonstrate tf32 in conv3d on MI30X platform MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform Sep 1, 2025
Comment thread projects/miopen/src/include/miopen/datatype.hpp Outdated
@yingluAMD yingluAMD force-pushed the conv_tf32_poc branch 2 times, most recently from e08d2bd to 2b0a177 Compare September 11, 2025 07:06
@yingluAMD yingluAMD force-pushed the conv_tf32_poc branch 2 times, most recently from a17b62f to 7e93a40 Compare September 15, 2025 08:45
@yingluAMD yingluAMD changed the title MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform [WIP] Sep 15, 2025
@yingluAMD yingluAMD changed the title MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform [WIP] MIOpen:feature:tf32:demonstrate tf32 in conv3d on MI30X platform Sep 15, 2025
Comment thread projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp Outdated
@yingluAMD yingluAMD force-pushed the conv_tf32_poc branch 2 times, most recently from 0da782e to 821e4f2 Compare September 18, 2025 02:32
@yingluAMD yingluAMD self-assigned this Sep 30, 2025
Copy link
Copy Markdown
Contributor

@BrianHarrisonAMD BrianHarrisonAMD left a comment

Choose a reason for hiding this comment

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

I believe this is ready to merge now.
You will need to update to latest from develop, and get TheRock CI passing before merge though.

Thanks for the work on this!

@yingluAMD yingluAMD merged commit 09ab715 into ROCm:develop Nov 16, 2025
31 checks passed
@yingluAMD yingluAMD mentioned this pull request Nov 18, 2025
1 task
lajagapp added a commit that referenced this pull request Nov 21, 2025
BradPepersAMD added a commit that referenced this pull request Nov 25, 2025
BradPepersAMD added a commit that referenced this pull request Nov 25, 2025
@yingluAMD yingluAMD mentioned this pull request Nov 26, 2025
1 task
@yingluAMD yingluAMD deleted the conv_tf32_poc branch April 1, 2026 03:47
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants