Skip to content

fix the int4 ark test #2038

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.

Sign up for GitHub

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

Open
a32543254 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

fix the int4 ark test #2038

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 4 additions & 8 deletions auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_fp8_dpas.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -733,14 +733,6 @@ CUTE_DEVICE void MoEGEMM(const ElementA* Activations, const ElementB* Weights,
int group_range = item.get_group_range(1);
int local_id = item.get_local_linear_id();

if (group_id == 0 && local_id == 0) {
auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::global_space>(
atomic_buffer[0]);
atm.store(0);
}

int pre_rows = 0;
int pre_tiles = 0;

Expand Down Expand Up @@ -935,6 +927,8 @@ void moe_prefill_fp8_dpas_per_tensor_dispatch(
throw std::runtime_error(
"moe_prefill_fp8_dpas(per-tensor): failed to allocate atomic buffer");
}
// Initialize the persistent scheduler counter before kernel launch.
q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();

#define ARK_DPAS_PT_LAUNCH(policy) \
MoEGEMMLauncher<'R', 'R', policy, ScaleMode::kPerTensor>( \
Expand Down Expand Up @@ -993,6 +987,8 @@ void moe_prefill_fp8_dpas_per_group_dispatch(
throw std::runtime_error(
"moe_prefill_fp8_dpas(per-group): failed to allocate atomic buffer");
}
// Initialize the persistent scheduler counter before kernel launch.
q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();

#define ARK_DPAS_PG_LAUNCH(policy) \
MoEGEMMLauncher<'R', 'C', policy, ScaleMode::kPerGroup>( \
Expand Down
12 changes: 4 additions & 8 deletions auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_int_dpas.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -584,14 +584,6 @@ CUTE_DEVICE void MoEGEMM_int(const ElementA* Activations,
int group_range = item.get_group_range(1);
int local_id = item.get_local_linear_id();

if (group_id == 0 && local_id == 0) {
auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::global_space>(
atomic_buffer[0]);
atm.store(0);
}

int pre_rows = 0;
int pre_tiles = 0;

Expand Down Expand Up @@ -785,6 +777,8 @@ void moe_prefill_int_dpas_per_tensor_dispatch(
throw std::runtime_error(
"moe_prefill_int_dpas(per-tensor): failed to allocate atomic buffer");
}
// Initialize the persistent scheduler counter before kernel launch.
q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();

#define ARK_DPAS_INT_PT_LAUNCH(policy) \
MoEGEMMLauncher_int<'R', 'R', policy, ScaleMode::kPerTensor>( \
Expand Down Expand Up @@ -848,6 +842,8 @@ void moe_prefill_int_dpas_per_group_dispatch(
throw std::runtime_error(
"moe_prefill_int_dpas(per-group): failed to allocate atomic buffer");
}
// Initialize the persistent scheduler counter before kernel launch.
q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();

#define ARK_DPAS_INT_PG_LAUNCH_SYM(policy) \
MoEGEMMLauncher_int<'R', 'C', policy, ScaleMode::kPerGroup>( \
Expand Down
12 changes: 4 additions & 8 deletions auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_s4_dpas.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -423,14 +423,6 @@ CUTE_DEVICE void MoEGEMM_s4(const ElementA* Activations,
int group_range = item.get_group_range(1);
int local_id = item.get_local_linear_id();

if (group_id == 0 && local_id == 0) {
auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
sycl::memory_scope::device,
sycl::access::address_space::global_space>(
atomic_buffer[0]);
atm.store(0);
}

int pre_rows = 0;
int pre_tiles = 0;

Expand Down Expand Up @@ -642,6 +634,10 @@ void moe_prefill_s4_dpas_per_group_dispatch(
throw std::runtime_error(
"moe_prefill_s4_dpas(per-group): failed to allocate atomic buffer");
}
// Initialize the persistent scheduler counter on the host before launch.
// Doing this in-kernel without a grid-wide barrier can race with early
// work-groups that already execute atomicAdd on the same counter.
q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();

#define ARK_DPAS_S4_PG_LAUNCH_SYM(policy) \
MoEGEMMLauncher_s4<'R', 'C', policy>( \
Expand Down
Loading