Skip to content
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

Only enable example 14 when the backend is CUDA #170

Merged

Conversation

FMarno
Copy link
Collaborator

@FMarno FMarno commented Dec 6, 2024

Maybe it was already broken, but the recent rebase for CUTLASS 3.6 introduced this line:

https://github.com/codeplaysoftware/cutlass-fork/blame/d49319f72cd9eeeb49830a94e23b3fdf8635fcc8/include/cutlass/gemm/device/gemm_universal_adapter.h#L461

Making use of DispatchPolicy::SubgroupSize. The example examples/14_ampere_tf32_tensorop_gemm/ampere_tf32_tensorop_gemm_cute.cu had set DispatchPolicy to

using DispatchPolicy = cutlass::gemm::MainloopSm80CpAsync<PipelineStages>;

leading to this compiler error:

/tmp/finlay/cutlass-fork/include/cutlass/gemm/device/gemm_universal_adapter.h:461:70: error: no member named 'SubgroupSize' in 'cutlass::gemm::MainloopSm80CpAsync<4>'
  461 |           kernel_properties{sycl_exp::sub_group_size<DispatchPolicy::SubgroupSize>}
      |                                                      ~~~~~~~~~~~~~~~~^

I have modified the cmake to only enable that example for nvidia devices.

Copy link
Collaborator

@aacostadiaz aacostadiaz left a comment

Choose a reason for hiding this comment

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

LGTM

@aacostadiaz aacostadiaz merged commit 0f258ae into codeplaysoftware:sycl-develop Dec 12, 2024
5 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.

3 participants