submission 99101

  ⋯ 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': 4, 'BLOCK_W': 256}, num_warps=8, num_stages=5),
          triton.Config({'BLOCK_H': 16, 'BLOCK_W': 64}, num_warps=8, num_stages=4),
+         triton.Config({'BLOCK_H': 8, 'BLOCK_W': 256}, num_warps=8, num_stages=4),
          
-         # L4 / средние размеры
+         # === 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),
+         triton.Config({'BLOCK_H': 8, 'BLOCK_W': 64}, num_warps=4, num_stages=4),
+         triton.Config({'BLOCK_H': 2, 'BLOCK_W': 256}, num_warps=4, num_stages=4),
          
-         # Маленькие тензоры / fallback
-         triton.Config({'BLOCK_H': 2, 'BLOCK_W': 64},  num_warps=4, num_stages=3),
+         # === Маленькие ===
+         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'],
  ⋯ 8 unchanged lines
      BLOCK_H: tl.constexpr, BLOCK_W: tl.constexpr
  ):
      """
-     Супер-оптимизированное 2D-блочное ядро Conv2D.
-     - 2D тайлинг (BLOCK_H x BLOCK_W) для максимальной локальности
-     - Оптимизированный порядок циклов: cin -> kh -> kw
-     - Минимум arithmetic в горячих циклах
-     - Правильное использование масок для граничных условий
+     ⚡ МАКСИМАЛЬНО ОПТИМИЗИРОВАННОЕ 2D-БЛОЧНОЕ ЯДРО CONV2D
+     
+     Ключевые оптимизации:
+     ✅ 2D-блокировка для максимальной L1 cache локальности
+     ✅ Порядок циклов: cin → kh → kw (оптимально для NCHW)
+     ✅ Pre-computed offsets вне горячих циклов
+     ✅ 2D broadcasting для эффективных операций
+     ✅ num_stages=5 для максимального prefetching
+     ✅ Правильные маски для граничных условий
+     ✅ += вместо acc = acc + для компиляторной оптимизации
      """
      
-     # === Декодирование Grid IDs ===
+     # === Grid Decoding ===
      pid_w = tl.program_id(0)
      pid_h = tl.program_id(1)
      pid_z = tl.program_id(2)
  ⋯ 1 unchanged lines
      batch_idx = pid_z // C_OUT
      out_ch = pid_z % C_OUT
      
-     # === 2D Offsets (блочная обработка) ===
-     # Height offsets: [BLOCK_H]
+     # === 2D Output Offsets ===
      offs_h = pid_h * BLOCK_H + tl.arange(0, BLOCK_H)
      mask_h = offs_h < H_OUT
      
-     # Width offsets: [BLOCK_W]
      offs_w = pid_w * BLOCK_W + tl.arange(0, BLOCK_W)
      mask_w = offs_w < W_OUT
      
-     # === Базовые указатели ===
-     # Output: [batch, out_ch, h, w]
+     # === Base Pointers ===
      out_base = output_ptr + batch_idx * stride_out_n + out_ch * stride_out_c
-     
-     # Input: [batch, ...]
      in_base = input_ptr + batch_idx * stride_in_n
-     
-     # Weight: [out_ch, ...]
      wei_base = weight_ptr + out_ch * stride_w_out
      
-     # === Аккумулятор (2D блок в регистрах) ===
-     # [BLOCK_H, BLOCK_W] - используем всю вычислительную мощь
+     # === 2D Accumulator ===
      acc = tl.zeros([BLOCK_H, BLOCK_W], dtype=tl.float32)
      
-     # === ОСНОВНОЙ ЦИКЛ СВЕРТКИ ===
-     # Порядок: cin -> kh -> kw (оптимально для NCHW памяти)
+     # === Combined 2D Mask (early compute) ===
+     mask_2d = (mask_h[:, None] & mask_w[None, :])
+     
+     # === Main Convolution Loop ===
      for cin in range(C_IN):
          in_ch = in_base + cin * stride_in_c
          wei_ch = wei_base + cin * stride_w_in
          
          for kh in range(K):
-             # 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 pointer
+             # Pre-compute height offsets for this kernel position
+             in_h_off = (offs_h[:, None] + kh) * stride_in_h
+             in_row = in_ch + in_h_off
              wei_row = wei_ch + kh * stride_w_h
              
              for kw in range(K):
-                 # === Загрузка Веса ===
-                 # Один скаляр, используется для всего блока (broadcast)
+                 # Load weight (scalar, broadcasts to [BLOCK_H, BLOCK_W])
                  wei_val = tl.load(wei_row + kw * stride_w_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
+                 # Load input block (2D) with proper indexing
+                 in_w_off = (offs_w[None, :] + kw) * stride_in_w
+                 in_ptrs = in_row + in_w_off
+                 in_val = tl.load(in_ptrs, mask=mask_2d, other=0.0)
                  
-                 # Маска: обе координаты должны быть валидны
-                 in_val = tl.load(
-                     in_ptrs, 
-                     mask=(mask_h[:, None] & mask_w[None, :]), 
-                     other=0.0
-                 )
-                 
-                 # === FMA ===
+                 # FMA with += for better compiler optimization
                  acc += in_val * wei_val
      
-     # === ЗАПИСЬ РЕЗУЛЬТАТА ===
-     # 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
+     # === Store Result ===
+     out_h_off = offs_h[:, None] * stride_out_h
+     out_w_off = offs_w[None, :] * stride_out_w
+     out_ptrs = out_base + out_h_off + out_w_off
      
-     # Запись с маской для граничных условий
-     tl.store(
-         out_ptrs, 
-         acc, 
-         mask=(mask_h[:, None] & mask_w[None, :])
-     )
+     tl.store(out_ptrs, acc, mask=mask_2d)
   
   
  def custom_kernel(data):
-     """Оптимальная wrapper для Conv2D."""
+     """Optimized wrapper function."""
      
      input_tensor, kernel, output_tensor = data
      
-     # Гарантируем контигуозность
+     # Ensure contiguous memory layout
      input_tensor = input_tensor.contiguous()
      kernel = kernel.contiguous()
      
-     # Размеры
+     # Extract dimensions
      batch, c_in, h_in, w_in = input_tensor.shape
      c_out, _, k_h, k_w = kernel.shape
      
+     # Calculate output dimensions (stride=1, padding=0)
      h_out = h_in - k_h + 1
      w_out = w_in - k_w + 1
      
-     # Grid: (W_blocks, H_blocks, N*C_out)
+     # Grid configuration: (W_blocks, H_blocks, Batch*Out_Channels)
      grid = lambda META: (
          triton.cdiv(w_out, META['BLOCK_W']),
          triton.cdiv(h_out, META['BLOCK_H']),
          batch * c_out
      )
      
+     # Launch kernel
      conv2d_kernel_tiled[grid](
          input_tensor, kernel, output_tensor,
          *input_tensor.stride(),