Replace std::min with cute::min in sm120 blockwise scaling device functions#3055
Open
vruga wants to merge 1 commit intoNVIDIA:mainfrom
Open
Replace std::min with cute::min in sm120 blockwise scaling device functions#3055vruga wants to merge 1 commit intoNVIDIA:mainfrom
vruga wants to merge 1 commit intoNVIDIA:mainfrom
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #3051.
Fix
Replaced
std::minwithcute::min, which is taggedCUTE_HOST_DEVICE constexprand is safe to call from device code. It accepts plain arithmetic types (both
arguments here are runtime integers), and is already the standard choice for this
pattern across the rest of the CUTLASS collective headers.
The same bug was also present in the
load()function of the sister filesm120_mma_array_tma_blockwise_scaling.hpp, so that has been fixed here as well.