deepseek-ai/DeepEP

Fork 0

mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00

[GH-ISSUE #346] question about half_warp_reduce_max #241

New issue

Closed

opened 2026-10-26 09:57:14 +00:00 by bsakaguchi · 4 comments

bsakaguchi commented

2026-10-26 09:57:14 +00:00

Owner

Originally created by @lpoyster on GitHub (Jul 31, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/346

I found that in the previous version, low latency fp8 cast use the following code, i wonder why the mask will be {0xffffffff, 0xffff} rather than random, thanks~

__forceinline__ __device__ float half_warp_reduce_max(float value) {
    auto mask = __activemask();
    // The mask be in `{0xffffffff, 0xffff}`
    value = max(value, __shfl_xor_sync(mask, value, 8));
    value = max(value, __shfl_xor_sync(mask, value, 4));
    value = max(value, __shfl_xor_sync(mask, value, 2));
    value = max(value, __shfl_xor_sync(mask, value, 1));
    return value;
}

Originally created by @lpoyster on GitHub (Jul 31, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/346 I found that in the previous version, low latency fp8 cast use the following code, i wonder why the mask will be {0xffffffff, 0xffff} rather than random, thanks~ ```cpp __forceinline__ __device__ float half_warp_reduce_max(float value) { auto mask = __activemask(); // The mask be in `{0xffffffff, 0xffff}` value = max(value, __shfl_xor_sync(mask, value, 8)); value = max(value, __shfl_xor_sync(mask, value, 4)); value = max(value, __shfl_xor_sync(mask, value, 2)); value = max(value, __shfl_xor_sync(mask, value, 1)); return value; } ```

bsakaguchi closed this issue

2026-10-26 09:57:14 +00:00

bsakaguchi commented

2026-10-26 09:57:15 +00:00

Author

Owner

@MARD1NO commented on GitHub (Aug 18, 2025):

I found that in the previous version, low latency fp8 cast use the following code, i wonder why the mask will be {0xffffffff, 0xffff} rather than random, thanks~

forceinline device float half_warp_reduce_max(float value) {
auto mask = __activemask();
// The mask be in {0xffffffff, 0xffff}
value = max(value, __shfl_xor_sync(mask, value, 8));
value = max(value, __shfl_xor_sync(mask, value, 4));
value = max(value, __shfl_xor_sync(mask, value, 2));
value = max(value, __shfl_xor_sync(mask, value, 1));
return value;
}

In stride loop, maybe only a half warp execute the reduce_max, so the mask should be 0xffff, otherwise, it will get deadlock if you use full mask

@MARD1NO commented on GitHub (Aug 18, 2025): > I found that in the previous version, low latency fp8 cast use the following code, i wonder why the mask will be {0xffffffff, 0xffff} rather than random, thanks~ > > __forceinline__ __device__ float half_warp_reduce_max(float value) { > auto mask = __activemask(); > // The mask be in `{0xffffffff, 0xffff}` > value = max(value, __shfl_xor_sync(mask, value, 8)); > value = max(value, __shfl_xor_sync(mask, value, 4)); > value = max(value, __shfl_xor_sync(mask, value, 2)); > value = max(value, __shfl_xor_sync(mask, value, 1)); > return value; > } In stride loop, maybe only a half warp execute the reduce_max, so the mask should be 0xffff, otherwise, it will get deadlock if you use full mask

bsakaguchi commented

2026-10-26 09:57:15 +00:00

Author

Owner

@lpoyster commented on GitHub (Aug 25, 2025):

thx for reply, i wonder why the mask will exactly be {0xffffffff, 0xffff}，for example, only one thread executing this instruction at some moment @MARD1NO

@lpoyster commented on GitHub (Aug 25, 2025): thx for reply, i wonder why the mask will exactly be {0xffffffff, 0xffff}，for example, only one thread executing this instruction at some moment @MARD1NO

bsakaguchi commented

2026-10-26 09:57:15 +00:00

Author

Owner

@MARD1NO commented on GitHub (Aug 25, 2025):

thx for reply, i wonder why the mask will exactly be {0xffffffff, 0xffff}，for example, only one thread executing this instruction at some moment @MARD1NO

Deepep use pertoken perblock 128 quantize, by using vectorized load, each thread can load 128bit = 8xbf16, and 16thread can process a 128 block, so the minimum threads here is 16

@MARD1NO commented on GitHub (Aug 25, 2025): > thx for reply, i wonder why the mask will exactly be {0xffffffff, 0xffff}，for example, only one thread executing this instruction at some moment [@MARD1NO](https://github.com/MARD1NO) Deepep use pertoken perblock 128 quantize, by using vectorized load, each thread can load 128bit = 8xbf16, and 16thread can process a 128 block, so the minimum threads here is 16

bsakaguchi commented

2026-10-26 09:57:15 +00:00

Author

Owner

@sphish commented on GitHub (Aug 25, 2025):

@lpoyster Sorry for the late reply. We previously had a misunderstanding about __activemask(). Your understanding is correct, and we have already fixed this issue in the latest code.

@sphish commented on GitHub (Aug 25, 2025): @lpoyster Sorry for the late reply. We previously had a misunderstanding about `__activemask()`. Your understanding is correct, and we have already fixed this issue in the latest code.

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepEP#241

No description provided.

Rows
Columns