Skip to content

Commit fbf6c16

Browse files
authored
[None][fix] Update the default invalid value for deepseek mode of routing (#9222)
Signed-off-by: Christina Zhang <[email protected]>
1 parent 92f86a5 commit fbf6c16

File tree

4 files changed

+8
-7
lines changed

4 files changed

+8
-7
lines changed

cpp/tensorrt_llm/kernels/noAuxTcKernels.cu

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
#include "tensorrt_llm/common/cudaTypeUtils.cuh"
2020
#include "tensorrt_llm/common/envUtils.h"
2121
#include "tensorrt_llm/kernels/noAuxTcKernels.h"
22+
#include <cmath>
2223
#include <cooperative_groups.h>
2324
#include <cooperative_groups/reduce.h>
2425

@@ -76,7 +77,7 @@ __global__ void deepseek_v3_topk_kernel(InputT* scores, OutputT* topkValues, Idx
7677
// note that for invalid scores, we simply use a negative value:
7778
// they work well even with the compacted format used in topK, and
7879
// sigmoid / bias activated scores cannot be negative
79-
static constexpr float invalidScoreFloat = -1.F;
80+
static constexpr float invalidScoreFloat = float{-INFINITY};
8081
const OutputT invalidScore = OutputT{invalidScoreFloat};
8182

8283
// load bias already; each warp represents one expert group

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ __global__ void routingMainKernel(KernelParams params)
6363
// note that for invalid scores, we simply use a negative value:
6464
// they work well even with the compacted format used in topK, and
6565
// sigmoid / bias activated scores cannot be negative
66-
static constexpr float invalidScoreFloat = -1.F;
66+
static constexpr float invalidScoreFloat = float{-INFINITY};
6767
const OutputT invalidScore = OutputT{invalidScoreFloat};
6868

6969
// load bias already; each warp represents one expert group

cpp/tests/unit_tests/kernels/routing/routingDeepSeekTest.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -49,7 +49,7 @@ class RoutingDeepSeekKernelTest : public RoutingKernelTest<T>
4949
// note that for invalid scores, we simply use a negative value:
5050
// they work well even with the compacted format used in topK, and
5151
// sigmoid / bias activated scores cannot be negative
52-
static constexpr float invalidScoreFloat = -1.F;
52+
static constexpr float invalidScoreFloat = float{-INFINITY};
5353
const T invalidScore = T{invalidScoreFloat};
5454

5555
float scoreSigmoid[param.numExperts];

cpp/tests/unit_tests/kernels/routing/routingTest.h

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -17,17 +17,17 @@
1717

1818
#include <gtest/gtest.h>
1919

20-
#include <chrono>
21-
#include <memory> //@todo check the usage of this
22-
#include <random> //@todo check the usage of this
23-
2420
#include "tensorrt_llm/common/memoryUtils.h"
2521
#include "tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.h"
2622
#include "tensorrt_llm/runtime/bufferManager.h"
2723
#include "tensorrt_llm/runtime/cudaStream.h"
2824
#include "tensorrt_llm/runtime/iBuffer.h"
2925
#include "tensorrt_llm/runtime/runtimeKernels.h"
3026
#include "tensorrt_llm/runtime/tllmLogger.h"
27+
#include <chrono>
28+
#include <cmath>
29+
#include <memory> //@todo check the usage of this
30+
#include <random> //@todo check the usage of this
3131

3232
namespace tensorrt_llm::tests::kernels::routing
3333
{

0 commit comments

Comments
 (0)