Skip to content

[Metax][CI]: skip trap asm on MetaX GPU to fix compile error #7975

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
zhang-chenyi wants to merge 1 commit into
base: develop
Choose a base branch
from
+2 −1
Open

[Metax][CI]: skip trap asm on MetaX GPU to fix compile error #7975

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
3 changes: 2 additions & 1 deletion custom_ops/gpu_ops/append_attn/mla_cache_kernel.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -212,14 +212,15 @@ __global__ void prefill_absorb_cache_kernel(
const uint32_t block_idx = block_table_now[ori_seq_id / block_size];
const uint32_t block_offset = ori_seq_id % block_size;

#ifndef PADDLE_WITH_CUSTOM_DEVICE_METAX_GPU
const int32_t block_idx1 = slot_mapping[token_idx] / block_size;
if (block_idx1 != block_idx) {
printf("block_idx1 %d != block_idx %d\n", block_idx1, block_idx);
printf("token_idx %d\n", token_idx);
printf("slot_mapping %d\n", slot_mapping[token_idx]);
asm volatile("trap;");
}

#endif
if (bias < nope_hidden_size) { // pe
const uint32_t inner_bias = bias;
const uint32_t hi = inner_bias / nope_size;
Expand Down
Loading