| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | import torch |
| 4 | from torch._inductor.runtime.triton_helpers import libdevice |
| 5 | |
| 6 | from vllm.triton_utils import tl, triton |
| 7 | |
| 8 | |
| 9 | @triton.jit |
| 10 | def _num_nans_kernel( |
| 11 | logits_ptr, |
| 12 | logits_stride, |
| 13 | num_nans_ptr, |
| 14 | vocab_size, |
| 15 | BLOCK_SIZE: tl.constexpr, |
| 16 | ): |
| 17 | req_idx = tl.program_id(0) |
| 18 | num_nans = 0 |
| 19 | for i in range(0, vocab_size, BLOCK_SIZE): |
| 20 | block = i + tl.arange(0, BLOCK_SIZE) |
| 21 | mask = block < vocab_size |
| 22 | logits = tl.load( |
| 23 | logits_ptr + req_idx * logits_stride + block, mask=mask, other=0 |
| 24 | ) |
| 25 | logits = logits.to(tl.float32) |
| 26 | is_nan = libdevice.isnan(logits).to(tl.int1) |
| 27 | num_nans += tl.sum(is_nan).to(tl.int32) |
| 28 | tl.store(num_nans_ptr + req_idx, num_nans) |
| 29 | |
| 30 | |
| 31 | def get_num_nans(logits: torch.Tensor) -> torch.Tensor: |
| 32 | num_reqs, vocab_size = logits.shape |
| 33 | BLOCK_SIZE = 8192 |
| 34 | num_nans = torch.empty(num_reqs, dtype=torch.int32, device=logits.device) |
| 35 | _num_nans_kernel[(num_reqs,)]( |
| 36 | logits, |
| 37 | logits.stride(0), |
| 38 | num_nans, |
| 39 | vocab_size, |
| 40 | BLOCK_SIZE=BLOCK_SIZE, |
| 41 | ) |
| 42 | return num_nans |