I found some clues

  

      github.com
  

  
    pytorch/pytorch/blob/95f4cd0ba99b04a82d711a47033e16a69f8f1b5e/aten/src/ATen/native/cuda/TensorTopK.cpp#L32-L52 2



      
      
          TORCH_IMPL_FUNC(topk_out_cuda)
            (const Tensor& self,
             int64_t k, int64_t dim, bool largest, bool sorted,
             const Tensor& values,
             const Tensor& indices) {
            TensorArg topK_arg{values, "topK", 1}, indices_arg{indices, "indices", 2}, input_arg{self, "self", 3};
            checkAllSameGPU(__func__, {topK_arg, indices_arg, input_arg});
          
          
  dim = at::maybe_wrap_dim(dim, self);
          
          
  if (should_use_sort(self, dim)) {
              topk_out_with_sort(self, k, dim, largest, values, indices);
              return;
            }
          
          
  // If k is 0 the result is an empty tensor, so we don't need to launch a kernel.
            if (k == 0) {
              return;
            }
          
          
  launch_gather_topk_kernel(self, k, dim, largest, sorted, values, indices);
      
    
  


  

  
    
    
  

  



  

      github.com
  

  
    pytorch/pytorch/blob/4262c8913c2bddb8d91565888b4871790301faba/aten/src/ATen/native/cuda/TensorTopK.cu#L92-L121 4



      
      
          for (IndexType i = threadIdx.x; i < numIterations; i += blockDim.x) {
            bool inRange = (i < inputSliceSize);
            T v =
              inRange ? doLdg(&inputSliceStart[i * inputWithinSliceStride]) : static_cast<T>(0);
            const auto convertedV = at::native::TopKTypeConfig<T>::convert(v);
            bool hasTopK;
            if (Order) {
              hasTopK = inRange && (convertedV > topKConverted);
            } else {
              hasTopK = inRange && (convertedV < topKConverted);
            }
          
          
  int index;
            int carry;
            at::cuda::exclusiveBinaryPrefixScan<int, true>(
                smem, hasTopK, &index, &carry, AddOp<int>());
          
          
  if (hasTopK) {
              int writeIndex = writeIndexStart + index;
              CUDA_KERNEL_ASSERT(writeIndex < outputSliceSize);
      
    
  

  This file has been truncated. show original