File tree Expand file tree Collapse file tree 2 files changed +9
-0
lines changed Expand file tree Collapse file tree 2 files changed +9
-0
lines changed Original file line number Diff line number Diff line change 25
25
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26
26
#include < float.h>
27
27
#include < algorithm>
28
+ #include < limits>
28
29
#include " ../include/ctc_prefix_decoder_host.h"
29
30
#include " ctc_fast_divmod.cuh"
30
31
#include " cub/cub.cuh"
@@ -440,7 +441,11 @@ __launch_bounds__(BLOCK_SIZE) void topk_reduce_and_copy_list_per_batch_kernel(
440
441
topk_values,
441
442
beam,
442
443
items_per_batch,
444
+ #if CUDART_VERSION >= 12090 // CUDA 12.9 and later
445
+ std::numeric_limits<float >::lowest (),
446
+ #else
443
447
cub::FpLimits<float >::Lowest (),
448
+ #endif
444
449
block_topk_fun,
445
450
set_key_value);
446
451
Original file line number Diff line number Diff line change @@ -94,7 +94,11 @@ __global__ void falign_cuda_step_kernel(
94
94
alphas_a[curIdxOffset][i] = result + logProbs_a[batchIndex][t][labelIdx];
95
95
threadMax = max (threadMax, alphas_a[curIdxOffset][i]);
96
96
}
97
+ #if CUDART_VERSION >= 12090 // CUDA 12.9 and later
98
+ scalar_t maxResult = BlockReduce (tempStorage).Reduce (threadMax, thrust::maximum<scalar_t >());
99
+ #else
97
100
scalar_t maxResult = BlockReduce (tempStorage).Reduce (threadMax, cub::Max ());
101
+ #endif
98
102
if (threadIdx .x == 0 ) {
99
103
maxValue = maxResult;
100
104
}
You can’t perform that action at this time.
0 commit comments