submission 233822

  ⋯ 2 unchanged lines
  from task import input_t, output_t
  
  # =============================================================================
- # Inline CUDA: B200 TMA + LDS.128 + 1024ピクセル/タイル版
+ # Inline CUDA: 最小限最適化版 (TMA + LDS.128)
  # =============================================================================
  # 目標: 598.336μs以下（B200 1位）
  #
- # v2_14からの最適化 (EXP-01):
- #   タイルサイズを512→1024ピクセルに拡大
- #   - TMA発行回数が半減（524,288 → 262,144）
- #   - ブロック起動オーバーヘッドが半減
- #   - スレッド数を128→256に拡大
+ # v2_15 (598.976μs) からの最適化:
+ #   1. 境界チェック削除: n_pixelsは1024の倍数なので不要
+ #   2. コンパイラフラグ強化: -Xptxas=-O3 追加
  #
- # ベース設計 (v2_14から継承):
- #   - TMA (cp.async.bulk) による非同期バルクコピー
- #   - blockIdx.x による通常グリッド起動（cudaMalloc排除）
- #   - ld.shared.v4.f32 (LDS.128) で共有メモリから一括ロード
- #   - st.global.v4.f32 (STG.128) による出力ベクトル化
- #   - 1スレッド4ピクセル処理でスループット最大化
- #
- # 共有メモリ使用量: 12KB/ブロック (B200: 228KB/SM, 余裕あり)
+ # 注意: インターリーブは逆効果だったため、v2_15の計算順序を維持
  # =============================================================================
  
  cuda_src = """
  #include <cuda_runtime.h>
  
  // =============================================================================
- // 定数定義 (EXP-01: 1024ピクセル/タイル)
+ // 定数定義
  // =============================================================================
- constexpr int PIXELS_PER_TILE   = 1024;      // 1024ピクセル/タイル (512→1024)
- constexpr int BYTES_PER_TILE    = 12288;     // 1024 × 3 × 4 = 12288バイト (16の倍数)
- constexpr int FLOATS_PER_TILE   = 3072;      // 1024 × 3 = 3072 floats
+ constexpr int PIXELS_PER_TILE   = 1024;
+ constexpr int BYTES_PER_TILE    = 12288;
+ constexpr int FLOATS_PER_TILE   = 3072;
+ constexpr int THREADS_PER_BLOCK = 256;
+ constexpr int PIXELS_PER_THREAD = 4;
  
- constexpr int THREADS_PER_BLOCK = 256;       // 256スレッド/ブロック (128→256)
- constexpr int PIXELS_PER_THREAD = 4;         // 4ピクセル/スレッド維持 (256 × 4 = 1024)
- 
  // =============================================================================
  // PTXヘルパー関数
  // =============================================================================
  
- // mbarrier初期化
  __device__ __forceinline__ void mbarrier_init(uint64_t* mbar, int arrival_count) {
      asm volatile(
          "mbarrier.init.shared.b64 [%0], %1;"
  ⋯ 1 unchanged lines
      );
  }
  
- // mbarrier: 期待するトランザクションバイト数を設定 + 到着
  __device__ __forceinline__ void mbarrier_arrive_expect_tx(uint64_t* mbar, int tx_bytes) {
      asm volatile(
          "mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"
  ⋯ 1 unchanged lines
      );
  }
  
- // mbarrier: フェーズビットで待機
  __device__ __forceinline__ void mbarrier_wait(uint64_t* mbar, int phase) {
      int done = 0;
      while (!done) {
  ⋯ 9 unchanged lines
      }
  }
  
- // cp.async.bulk: グローバル→共有メモリのバルクコピー
  __device__ __forceinline__ void cp_async_bulk(
      void* smem_ptr,
      const void* gmem_ptr,
  ⋯ 9 unchanged lines
      );
  }
  
- // 共有メモリからのベクトルロード (LDS.128)
  __device__ __forceinline__ void load_shared_float4(
      uint32_t smem_addr,
      float& v0, float& v1, float& v2, float& v3
  ⋯ 5 unchanged lines
      );
  }
  
- // float4ベクトルストア（STG.128）
  __device__ __forceinline__ void store_float4_ptx(
      float* ptr, float v0, float v1, float v2, float v3
  ) {
  ⋯ 4 unchanged lines
  }
  
  // =============================================================================
- // TMA + LDS.128 + 1024ピクセル/タイル カーネル
+ // 最小限最適化カーネル（v2_15の構造維持、境界チェックのみ削除）
  // =============================================================================
- // __launch_bounds__(256, 4): SM当たり4ブロック（共有メモリ12KB×4=48KB < 228KB）
- __global__ __launch_bounds__(256, 4) void grayscale_1024tile_kernel(
+ __global__ __launch_bounds__(256, 4) void grayscale_minimal_opt_kernel(
      const float* __restrict__ input,
-     float* __restrict__ output,
-     const int n_pixels
+     float* __restrict__ output
  ) {
-     // 共有メモリ: 12KB (3072 floats) + mbarrier (8B)
      __shared__ __align__(128) float smem[FLOATS_PER_TILE];
      __shared__ __align__(8) uint64_t mbar;
  
  ⋯ 1 unchanged lines
      const int tile_idx = blockIdx.x;
      const int pixel_base = tile_idx * PIXELS_PER_TILE + tid * PIXELS_PER_THREAD;
  
-     // 境界チェック（268,435,456 / 1024 = 262,144 で割り切れるが、安全のため維持）
-     if (pixel_base >= n_pixels) return;
+     // 【最適化】境界チェック削除（n_pixelsは1024の倍数）
  
-     // ========================================
-     // Step 1: mbarrier初期化（Thread 0のみ）
-     // ========================================
      if (tid == 0) {
          mbarrier_init(&mbar, 1);
      }
-     __syncthreads();  // mbarrier初期化の可視性を保証
+     __syncthreads();
  
-     // ========================================
-     // Step 2: TMAロード発行（Thread 0のみ）
-     // ========================================
      if (tid == 0) {
          const float* src = input + tile_idx * PIXELS_PER_TILE * 3;
          mbarrier_arrive_expect_tx(&mbar, BYTES_PER_TILE);
          cp_async_bulk(smem, src, BYTES_PER_TILE, &mbar);
      }
  
-     // ========================================
-     // Step 3: TMA完了待機
-     // ========================================
      mbarrier_wait(&mbar, 0);
  
-     // ========================================
-     // Step 4: LDS.128で共有メモリから一括ロード
-     // ========================================
-     // tid * 12 floats = tid * 48 bytes (常に16の倍数)
+     // v2_15と同じ順序を維持（コンパイラに最適化を任せる）
      uint32_t smem_addr = (uint32_t)__cvta_generic_to_shared(&smem[tid * 12]);
  
-     // s0 = [R0, G0, B0, R1]
-     // s1 = [G1, B1, R2, G2]
-     // s2 = [B2, R3, G3, B3]
      float s0_x, s0_y, s0_z, s0_w;
      float s1_x, s1_y, s1_z, s1_w;
      float s2_x, s2_y, s2_z, s2_w;
  
      load_shared_float4(smem_addr,      s0_x, s0_y, s0_z, s0_w);
-     load_shared_float4(smem_addr + 16, s1_x, s1_y, s1_z, s1_w);  // +16バイト = +4floats
-     load_shared_float4(smem_addr + 32, s2_x, s2_y, s2_z, s2_w);  // +32バイト = +8floats
+     load_shared_float4(smem_addr + 16, s1_x, s1_y, s1_z, s1_w);
+     load_shared_float4(smem_addr + 32, s2_x, s2_y, s2_z, s2_w);
  
-     // ========================================
-     // Step 5: Grayscale計算 (fmafチェーン)
-     // ========================================
-     // マッピング:
-     // Pixel 0: R=s0_x, G=s0_y, B=s0_z
-     // Pixel 1: R=s0_w, G=s1_x, B=s1_y
-     // Pixel 2: R=s1_z, G=s1_w, B=s2_x
-     // Pixel 3: R=s2_y, G=s2_z, B=s2_w
      float gray0 = fmaf(s0_x, 0.2989f, fmaf(s0_y, 0.5870f, s0_z * 0.1140f));
      float gray1 = fmaf(s0_w, 0.2989f, fmaf(s1_x, 0.5870f, s1_y * 0.1140f));
      float gray2 = fmaf(s1_z, 0.2989f, fmaf(s1_w, 0.5870f, s2_x * 0.1140f));
      float gray3 = fmaf(s2_y, 0.2989f, fmaf(s2_z, 0.5870f, s2_w * 0.1140f));
  
-     // ========================================
-     // Step 6: STG.128でグローバルメモリに出力
-     // ========================================
      store_float4_ptx(output + pixel_base, gray0, gray1, gray2, gray3);
  }
  
- // =============================================================================
- // ホストインターフェース
- // =============================================================================
  torch::Tensor rgb_to_grayscale(torch::Tensor input, torch::Tensor output) {
      int n_pixels = output.numel();
+     int blocks = n_pixels / PIXELS_PER_TILE;
  
-     // ブロック数計算: 268,435,456 / 1024 = 262,144 ブロック
-     int blocks = (n_pixels + PIXELS_PER_TILE - 1) / PIXELS_PER_TILE;
- 
-     grayscale_1024tile_kernel<<<blocks, THREADS_PER_BLOCK>>>(
+     grayscale_minimal_opt_kernel<<<blocks, THREADS_PER_BLOCK>>>(
          input.data_ptr<float>(),
-         output.data_ptr<float>(),
-         n_pixels
+         output.data_ptr<float>()
      );
  
      return output;
  ⋯ 12 unchanged lines
      global _module
      if _module is None:
          _module = load_inline(
-             name='grayscale_cuda_1024tile',
+             name='grayscale_cuda_minimal_opt',
              cuda_sources=[cuda_src],
              cpp_sources=[cpp_src],
              functions=['rgb_to_grayscale'],
  ⋯ 2 unchanged lines
                  '--use_fast_math',
                  '-std=c++17',
                  '-arch=sm_90',
+                 '-Xptxas=-O3',
              ],
              verbose=False
          )