submission 99100

  ⋯ 3 unchanged lines
   
  @triton.autotune(
      configs=[
-         # === A100/H100/B200 (максимальная производительность) ===
+         # A100/H100/B200 - максимальная производительность
          triton.Config({'BLOCK_H': 8, 'BLOCK_W': 128}, num_warps=8, num_stages=5),
-         triton.Config({'BLOCK_H': 16, 'BLOCK_W': 64}, num_warps=8, num_stages=5),
          triton.Config({'BLOCK_H': 4, 'BLOCK_W': 256}, num_warps=8, num_stages=5),
+         triton.Config({'BLOCK_H': 16, 'BLOCK_W': 64}, num_warps=8, num_stages=4),
+         
+         # L4 / средние размеры
          triton.Config({'BLOCK_H': 4, 'BLOCK_W': 128}, num_warps=8, num_stages=4),
+         triton.Config({'BLOCK_H': 8, 'BLOCK_W': 64},  num_warps=4, num_stages=4),
          
-         # === L4 / средние размеры ===
-         triton.Config({'BLOCK_H': 4, 'BLOCK_W': 64}, num_warps=8, num_stages=4),
-         triton.Config({'BLOCK_H': 2, 'BLOCK_W': 128}, num_warps=4, num_stages=4),
-         triton.Config({'BLOCK_H': 8, 'BLOCK_W': 32}, num_warps=4, num_stages=4),
-         
-         # === Маленькие тензоры / fallback ===
-         triton.Config({'BLOCK_H': 1, 'BLOCK_W': 64}, num_warps=2, num_stages=3),
-         triton.Config({'BLOCK_H': 2, 'BLOCK_W': 32}, num_warps=2, num_stages=3),
+         # Маленькие тензоры / fallback
+         triton.Config({'BLOCK_H': 2, 'BLOCK_W': 64},  num_warps=4, num_stages=3),
+         triton.Config({'BLOCK_H': 1, 'BLOCK_W': 128}, num_warps=2, num_stages=3),
      ],
      key=['w_out', 'h_out', 'c_in', 'k_size'],
  )
  @triton.jit
- def conv2d_kernel_optimized(
+ def conv2d_kernel_tiled(
      input_ptr, weight_ptr, output_ptr,
      stride_in_n, stride_in_c, stride_in_h, stride_in_w,
      stride_w_out, stride_w_in, stride_w_h, stride_w_w,
  ⋯ 2 unchanged lines
      BLOCK_H: tl.constexpr, BLOCK_W: tl.constexpr
  ):
      """
-     ✅ МАКСИМАЛЬНО ОПТИМИЗИРОВАННОЕ ЯДРО CONV2D
-     
-     Оптимизации:
-     1. 2D-блокировка (BLOCK_H × BLOCK_W) для максимальной локальности
-     2. Правильный порядок циклов: cin → kh → kw
-     3. Pre-computed pointers для минимума arithmetic
-     4. 2D broadcasting для эффективных загрузок/записей
-     5. num_stages=5 для максимального prefetching
-     6. Правильные маски для граничных условий
+     Супер-оптимизированное 2D-блочное ядро Conv2D.
+     - 2D тайлинг (BLOCK_H x BLOCK_W) для максимальной локальности
+     - Оптимизированный порядок циклов: cin -> kh -> kw
+     - Минимум arithmetic в горячих циклах
+     - Правильное использование масок для граничных условий
      """
      
-     # === Grid IDs ===
+     # === Декодирование Grid IDs ===
      pid_w = tl.program_id(0)
      pid_h = tl.program_id(1)
      pid_z = tl.program_id(2)
      
-     # Декодирование Batch и Output Channel
      batch_idx = pid_z // C_OUT
      out_ch = pid_z % C_OUT
      
-     # === 2D Offsets ===
-     # Height: [BLOCK_H] -> reshape в [BLOCK_H, 1] для broadcasting
+     # === 2D Offsets (блочная обработка) ===
+     # Height offsets: [BLOCK_H]
      offs_h = pid_h * BLOCK_H + tl.arange(0, BLOCK_H)
      mask_h = offs_h < H_OUT
      
-     # Width: [BLOCK_W] -> reshape в [1, BLOCK_W] для broadcasting
+     # Width offsets: [BLOCK_W]
      offs_w = pid_w * BLOCK_W + tl.arange(0, BLOCK_W)
      mask_w = offs_w < W_OUT
      
-     # === Базовые Pointers ===
-     # Output: [batch, out_ch, :, :]
+     # === Базовые указатели ===
+     # Output: [batch, out_ch, h, w]
      out_base = output_ptr + batch_idx * stride_out_n + out_ch * stride_out_c
      
-     # Input: [batch, :, :, :]
+     # Input: [batch, ...]
      in_base = input_ptr + batch_idx * stride_in_n
      
-     # Weight: [out_ch, :, :, :]
+     # Weight: [out_ch, ...]
      wei_base = weight_ptr + out_ch * stride_w_out
      
-     # === 2D Accumulator в Register File ===
-     # [BLOCK_H, BLOCK_W] - максимум параллелизма
+     # === Аккумулятор (2D блок в регистрах) ===
+     # [BLOCK_H, BLOCK_W] - используем всю вычислительную мощь
      acc = tl.zeros([BLOCK_H, BLOCK_W], dtype=tl.float32)
      
      # === ОСНОВНОЙ ЦИКЛ СВЕРТКИ ===
-     # Порядок: cin → kh → kw (оптимально для памяти NCHW)
+     # Порядок: cin -> kh -> kw (оптимально для NCHW памяти)
      for cin in range(C_IN):
-         # Pre-compute channel-shifted pointers
          in_ch = in_base + cin * stride_in_c
          wei_ch = wei_base + cin * stride_w_in
          
          for kh in range(K):
-             # Pre-compute Height offset (Broadcasting: [BLOCK_H, 1])
-             # offs_h[:, None] превращает [BLOCK_H] в [BLOCK_H, 1]
-             in_h = (offs_h[:, None] + kh) * stride_in_h
-             in_row = in_ch + in_h
+             # Pre-compute Input row pointer для этого kh
+             # offs_h[:, None] даёт размер [BLOCK_H, 1]
+             # Broadcasting: (BLOCK_H, 1) + скаляр = (BLOCK_H, 1)
+             in_h_idx = (offs_h[:, None] + kh) * stride_in_h
+             in_row = in_ch + in_h_idx
              
-             # Pre-compute Weight row offset
+             # Pre-compute Weight row pointer
              wei_row = wei_ch + kh * stride_w_h
              
              for kw in range(K):
-                 # === Загрузка Веса (Scalar broadcast) ===
-                 # Один скаляр используется для всего блока [BLOCK_H, BLOCK_W]
+                 # === Загрузка Веса ===
+                 # Один скаляр, используется для всего блока (broadcast)
                  wei_val = tl.load(wei_row + kw * stride_w_w)
                  
-                 # === Загрузка Входа (2D Block) ===
-                 # offs_w[None, :] превращает [BLOCK_W] в [1, BLOCK_W]
-                 # Broadcasting: [BLOCK_H, 1] + [1, BLOCK_W] = [BLOCK_H, BLOCK_W]
-                 in_w = (offs_w[None, :] + kw) * stride_in_w
-                 in_ptrs = in_row + in_w
+                 # === Загрузка Входа (2D блок) ===
+                 # offs_w[None, :] даёт размер [1, BLOCK_W]
+                 # Broadcasting: (BLOCK_H, 1) + [1, BLOCK_W] = (BLOCK_H, BLOCK_W)
+                 in_w_idx = (offs_w[None, :] + kw) * stride_in_w
+                 in_ptrs = in_row + in_w_idx
                  
-                 # Загрузка с правильной маской
+                 # Маска: обе координаты должны быть валидны
                  in_val = tl.load(
-                     in_ptrs,
-                     mask=(mask_h[:, None] & mask_w[None, :]),
+                     in_ptrs, 
+                     mask=(mask_h[:, None] & mask_w[None, :]), 
                      other=0.0
                  )
                  
-                 # === FMA (Fused Multiply-Add) ===
+                 # === FMA ===
                  acc += in_val * wei_val
      
      # === ЗАПИСЬ РЕЗУЛЬТАТА ===
-     # 2D адреса: base + h_offset * stride_h + w_offset * stride_w
-     out_h = offs_h[:, None] * stride_out_h
-     out_w = offs_w[None, :] * stride_out_w
-     out_ptrs = out_base + out_h + out_w
+     # Output addresses: base + h_offset * stride_h + w_offset * stride_w
+     out_h_idx = offs_h[:, None] * stride_out_h
+     out_w_idx = offs_w[None, :] * stride_out_w
+     out_ptrs = out_base + out_h_idx + out_w_idx
      
      # Запись с маской для граничных условий
      tl.store(
-         out_ptrs,
-         acc,
+         out_ptrs, 
+         acc, 
          mask=(mask_h[:, None] & mask_w[None, :])
      )
   
   
  def custom_kernel(data):
-     """✅ Оптимальная wrapper для Conv2D."""
+     """Оптимальная wrapper для Conv2D."""
      
      input_tensor, kernel, output_tensor = data
      
-     # Гарантируем контигуозность памяти
+     # Гарантируем контигуозность
      input_tensor = input_tensor.contiguous()
      kernel = kernel.contiguous()
      
  ⋯ 4 unchanged lines
      h_out = h_in - k_h + 1
      w_out = w_in - k_w + 1
      
-     # Grid: (W_blocks, H_blocks, Batch*C_out)
+     # Grid: (W_blocks, H_blocks, N*C_out)
      grid = lambda META: (
          triton.cdiv(w_out, META['BLOCK_W']),
          triton.cdiv(h_out, META['BLOCK_H']),
          batch * c_out
      )
      
-     # Запуск ядра
-     conv2d_kernel_optimized[grid](
+     conv2d_kernel_tiled[grid](
          input_tensor, kernel, output_tensor,
          *input_tensor.stride(),
          *kernel.stride(),