Skip to content

FMHA examples: use cute::min in device functions#3222

Open
LwhJesse wants to merge 1 commit into
NVIDIA:mainfrom
LwhJesse:fmha-use-cute-min
Open

FMHA examples: use cute::min in device functions#3222
LwhJesse wants to merge 1 commit into
NVIDIA:mainfrom
LwhJesse:fmha-use-cute-min

Conversation

@LwhJesse
Copy link
Copy Markdown

Use cute::min instead of std::min in FMHA device functions.

These helpers are marked CUTLASS_DEVICE, and cute::min is already the CuTe/CUTLASS host-device constexpr utility used for this pattern.

This matches the existing CUTLASS/CuTe device-code style and avoids relying on std::min in device functions.

Testing:

cmake -S . -B build \
  -DCMAKE_BUILD_TYPE=Release \
  -DCUTLASS_NVCC_ARCHS=90a \
  -DCUTLASS_ENABLE_TESTS=OFF \
  -DCUTLASS_ENABLE_EXAMPLES=ON

cmake --build build --target 88_hopper_fmha -j$(nproc)

Result:

Built target 88_hopper_fmha

I did not build the Blackwell FMHA target locally.

@LwhJesse
Copy link
Copy Markdown
Author

Hi maintainers, gentle ping on this small FMHA examples cleanup.

This replaces std::min with cute::min inside CUTLASS_DEVICE helpers, matching the existing CuTe/CUTLASS device-code style. I verified that 88_hopper_fmha builds locally, but I do not have local Blackwell FMHA coverage, so maintainer CI/review would be helpful.

Is there a preferred owner for this examples/FMHA area?

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.

1 participant