ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-19 22:39:03 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#4819 (commit b995a0b)

Browse Source 
[CK] Fix windows build issues

## Motivation

Full build on Windows is currently broken due to compiler errors, this
PR should help fix that. This is also holding up the following PR in the
TheRock: https://github.com/ROCm/TheRock/pull/3382

## Technical Details

1. I don't see a good reason to be nesting a windows include inside the
ck_tile namespace. It was causing compiler errors too: Windows.h comes
with min and max, which was conflicting with ck_tile::min and
ck_tile::max, so I moved it out. I also defined NOMINMAX to prevent this
inclusion in the future.
2. The TRUE/FALSE macros are already used by Windows.h, which causes an
error. So I've opted for True/False. You can see this pattern in other
rocm-libraries.
3. The M_PI macro isn't available, at least in the WIN32_LEAN_AND_MEAN
context, from \<cmath\> on Windows. We'll be able to use
std::numbers::v_pi\<float\> when we have C++20 support.
4. There was a missing \<chrono\> include.

## Test Plan

Test locally and make sure this doesn't impact existing CI.

## Test Result

Compiles locally and passes existing ci.

## Submission Checklist

- [ x ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This commit is contained in:
Brock Hargreaves
2026-02-25 16:13:13 +00:00
committed by assistant-librarian[bot]
parent a32d704d89
commit abf13bdec1
4 changed files with 25 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_ex.hpp
View File

@@ -109,8 +109,8 @@ struct FusedMoeGemmPipeline_FlatmmEx
constexpr auto NEG1 = number<-1>{};
constexpr auto I0 = number<0>{};
constexpr auto I1 = number<1>{};
constexpr auto TRUE = bool_constant<true>{};
constexpr auto FALSE = bool_constant<false>{};
constexpr auto True = bool_constant<true>{};
constexpr auto False = bool_constant<false>{};
CK_TILE_LDS_ADDR ADataType* smem_0 = reinterpret_cast<CK_TILE_LDS_ADDR ADataType*>(smem);
CK_TILE_LDS_ADDR ADataType* smem_1 = reinterpret_cast<CK_TILE_LDS_ADDR ADataType*>(
@@ -291,7 +291,7 @@ struct FusedMoeGemmPipeline_FlatmmEx
g_win.bottom_tensor_view_ = u_view;
}
}
load_tile_raw(g_, g_win, i_access, FALSE, PreNop{});
load_tile_raw(g_, g_win, i_access, False, PreNop{});
};
auto move_g = [&]() {
move_tile_window(g_win, {number<0>{}, number<BlockShape::Block_Kr0>{}, number<0>{}});
@@ -300,7 +300,7 @@ struct FusedMoeGemmPipeline_FlatmmEx
auto gld_d =
[&]<typename PreNop = bool_constant<false>>(auto& d_, auto i_access, PreNop = {}) {
load_tile_raw(d_, d_win, i_access, FALSE, PreNop{});
load_tile_raw(d_, d_win, i_access, False, PreNop{});
};
auto move_d = [&]() {
// d move along gemm-n
@@ -309,7 +309,7 @@ struct FusedMoeGemmPipeline_FlatmmEx
auto atomic_add_o =
[&]<typename PreNop = bool_constant<false>>(auto& o_, auto i_access, PreNop = {}) {
update_tile_raw(o_win, o_, i_access, TRUE, PreNop{});
update_tile_raw(o_win, o_, i_access, True, PreNop{});
};
auto acc_0 = Policy::template MakeCBlockTile_Gemm0<Problem>();
@@ -502,9 +502,9 @@ struct FusedMoeGemmPipeline_FlatmmEx
static_for<0, total_loops, 1>{}([&](auto i_issue) {
constexpr auto last_nop = [&]() {
if constexpr(i_issue == (total_loops - 1))
return TRUE;
return True;
else
return FALSE;
return False;
}();
gemm_0(acc_0, as[I1], gs[I1], i_issue, last_nop); // last gemm has nop
});
@@ -607,8 +607,8 @@ struct FusedMoeGemmPipeline_FlatmmEx
// start of pipeline
// clang-format off
gld_a(a_sst_win0, NEG1, TRUE);
gld_g(gs[I0], NEG1, TRUE);
gld_a(a_sst_win0, NEG1, True);
gld_g(gs[I0], NEG1, True);
move_a();
move_g();
clear_tile(acc_0);