Skip to content

Commit d8c5dfb

Browse files
committed
issue/899 - fix: fix causal_softmax and rearrange bug
1 parent 3c8fb3c commit d8c5dfb

3 files changed

Lines changed: 114 additions & 244 deletions

File tree

src/infiniop/ops/causal_softmax/moore/causal_softmax_kernel.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@ __device__ void causalSoftmaxKernel(
2828
// 1 | * * * ... * * |
2929
// 2 | * * * ... * * * |
3030
// height: 3 col_id->
31-
if (width + blockIdx.x >= threadIdx.x + height) {
31+
if (width + blockIdx.x >= col + height) {
3232
if constexpr (std::is_same_v<Tdata, half> || std::is_same_v<Tdata, cuda_bfloat16>) {
3333
/*
3434
* MUSA does not support CUDA's native `hexp` function.

0 commit comments

Comments
 (0)