topk unreachable code

[kurquhar]

Describe the issue

The following code is unreachable:

      } else {  // reorder by ascending index
        ExcludeOutput<int64_t><<<blocks_per_grid_D, BT, 0, stream>>>(output_value, K, dimension);
        CUDA_RETURN_IF_ERROR(cub::DeviceRadixSort::SortPairs(temp_storage, temp_bytes, output_value, input_value, output_key, input_key, dimension, 0, sizeof(T) * 8, stream));
        FillOutput<CudaT><<<blocks_per_grid_K, BT, 0, stream>>>(input_key, input_value, output_v_ptr, output_i, elem_nums, size, axis, K, i, dimension);
      }

https://github.com/microsoft/onnxruntime/blob/main/onnxruntime/core/providers/cuda/math/topk_impl.cuh#L463

due to this previous (0 == sorted) code path:
} else if (K <= BT * 16 || 0 == sorted) {
https://github.com/microsoft/onnxruntime/blob/main/onnxruntime/core/providers/cuda/math/topk_impl.cuh#L416

To reproduce

code inspection

Urgency

Unclear if the unreachable code is needed for correct behavior. Reference
It seems to be sorting by index to meet this statement (present since v1):
Given two equivalent values, this operator uses the indices along the axis as a tiebreaker. That is, the element with the lower index will appear first.
But, the same spec says that unsorted order is undefined (new in v11):
If “sorted” is 0, order of returned ‘Values’ and ‘Indices’ are undefined.

Platform

Linux

OS Version

Ubuntu 22.04

ONNX Runtime Installation

Built from Source

ONNX Runtime Version or Commit ID

15cb2f5

ONNX Runtime API

C++

Architecture

X64

Execution Provider

CUDA

Execution Provider Library Version

No response

[pranavsharma]

Let us know if this is causing an issue in your application. Thanks!