submission 673544

  # -*- coding: utf-8 -*-
  import os
- import logging
  import torch
  
- # Determinismo de ejecución en operaciones atómicas y de ordenamiento
+ # 1. ESTABILIZACIÓN DE LA WORLD LINE (Fuera del hot-path)
+ # Seteamos el determinismo globalmente al cargar el módulo para ahorrar microsegundos
  os.environ["CUBLAS_WORKSPACE_CONFIG"] = ":4096:8"
+ torch.use_deterministic_algorithms(True)
  
- # Configuración de telemetría obligatoria
- logging.basicConfig(level=logging.INFO)
- logger = logging.getLogger("Sentinel_Omega_Sort")
+ # 2. CACHÉ DE VÉRTICES (Persistent Index Buffer)
+ # Pre-asignamos un buffer de índices para evitar que torch.sort 
+ # tenga que pedir memoria nueva al driver de la GPU en cada iteración.
+ _INDEX_BUFFER = {}
  
  def custom_kernel(data):
      """
-     Motor de ordenamiento ascendente para arreglos 1D masivos.
-     Optimización basada en la evasión de asignación de memoria intermedia
-     y delegación al backend rocPRIM subyacente de PyTorch.
+     Sentinel Omega - Ordenamiento Radix v3.
+     Estrategia: Static Index Dispatch + Output In-Place.
      """
-     try:
-         # Desempaquetado estricto
-         input_tensor = data[0]
-         output = data[1]
-         
-         torch.use_deterministic_algorithms(True)
-         
-         # Ejecución del ordenamiento.
-         # torch.sort invoca rocprim::radix_sort en hardware AMD.
-         # Se captura exclusivamente el tensor de valores, omitiendo los índices.
-         sorted_vals, _ = torch.sort(input_tensor)
-         
-         # Copia directa en memoria (in-place) sobre el búfer de salida preasignado.
-         # Elimina la sobrecarga de instanciación de nuevos tensores en VRAM.
-         output.copy_(sorted_vals)
-         
-         return output
+     # 3. UNPACKING POR POSICIÓN (Fricción Cero)
+     # Evita el overhead de la gestión de excepciones de Python
+     input_t = data[0]
+     output_t = data[1]
+     
+     # 4. GESTIÓN DE LA CAJA DE ÍNDICES
+     # Si el tamaño cambia, el buffer se ajusta, pero en ráfagas de producción
+     # esto ocurre una sola vez, logrando latencia constante.
+     n = input_t.shape[0]
+     if n not in _INDEX_BUFFER:
+         # Los índices no los usaremos, pero torch.sort los requiere
+         _INDEX_BUFFER[n] = torch.empty((n,), dtype=torch.int64, device=input_t.device)
  
-     except Exception as e:
-         logger.error(f"Falla técnica en motor de ordenamiento: {str(e)}")
-         raise e
+     # 5. EL COLAPSO (Radix Sort Directo)
+     # Al usar el parámetro 'out', le decimos a rocPRIM que inyecte 
+     # los valores directamente en el buffer de salida (data[1]).
+     # Esto elimina 1 copia completa de memoria y 1 asignación (allocation).
+     torch.sort(input_t, out=(output_t, _INDEX_BUFFER[n]))
+     
+     return output_t