Skip to content

Commit ef47a10

Browse files
committed
fix(metax): skip trap asm on MetaX GPU to fix compile error
1 parent 91ca3d1 commit ef47a10

1 file changed

Lines changed: 2 additions & 0 deletions

File tree

custom_ops/gpu_ops/append_attn/mla_cache_kernel.cuh

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -212,13 +212,15 @@ __global__ void prefill_absorb_cache_kernel(
212212
const uint32_t block_idx = block_table_now[ori_seq_id / block_size];
213213
const uint32_t block_offset = ori_seq_id % block_size;
214214

215+
#ifndef PADDLE_WITH_CUSTOM_DEVICE_METAX_GPU
215216
const int32_t block_idx1 = slot_mapping[token_idx] / block_size;
216217
if (block_idx1 != block_idx) {
217218
printf("block_idx1 %d != block_idx %d\n", block_idx1, block_idx);
218219
printf("token_idx %d\n", token_idx);
219220
printf("slot_mapping %d\n", slot_mapping[token_idx]);
220221
asm volatile("trap;");
221222
}
223+
#endif
222224

223225
if (bias < nope_hidden_size) { // pe
224226
const uint32_t inner_bias = bias;

0 commit comments

Comments
 (0)