メニュー
Optimizes global, shared/LDS, and cache behavior for GPU kernels: coalescing, pipelining, bank conflicts, vector loads, and register vs shared tradeoffs. Use when NCU/rocprof shows low memory throughput, poor coalescing, or L2/LDS pressure.
Codex または Claude でインストール この Prompt をコピーして Codex、Claude、または他のアシスタントに貼り付けると、Skill ページを確認してインストールできます。
SOC 職業分類に基づく
Optimize fused cross-entropy loss kernels in Triton for NVIDIA and AMD GPUs. Covers fused log-softmax + NLL, online log-sum-exp, and large vocabulary handling. Use when writing or optimizing cross-entropy, label-smoothed CE, or similar loss kernels.
Optimize FlashAttention-style fused attention kernels in Triton for NVIDIA and AMD GPUs. Covers online softmax, tiled QK/AV GEMM, causal masking, and memory-efficient attention. Use when writing or optimizing self-attention, cross-attention, or any QKV attention kernel.
Optimize dense matrix multiplication (GEMM) kernels in Triton for NVIDIA and AMD GPUs. Covers tiled blocking, L2 cache grouping, tensor core utilization, and dual-platform autotune. Use when writing or optimizing matmul, linear layers, or any GEMM-based kernel.
Optimize RMS Normalization kernels in Triton for NVIDIA and AMD GPUs. Covers row-parallel reduction, vectorized loads, warp shuffle patterns, and rsqrt usage. Use when writing or optimizing RMSNorm, LayerNorm, or similar normalization kernels.
Optimize Rotary Position Embedding (RoPE) kernels in Triton for NVIDIA and AMD GPUs. Covers interleaved/sequential layouts, sincos precomputation, and multi-row processing. Use when writing or optimizing RoPE, position embeddings, or rotary transformations.
Optimize fused softmax kernels in Triton for NVIDIA and AMD GPUs. Covers online max/sum reduction, multi-row processing, numerical stability, and warp-level patterns. Use when writing or optimizing softmax, log-softmax, or similar row-wise normalization kernels.
| name | gpu-opt-tier2-memory-hierarchy |
| description | Optimizes global, shared/LDS, and cache behavior for GPU kernels: coalescing, pipelining, bank conflicts, vector loads, and register vs shared tradeoffs. Use when NCU/rocprof shows low memory throughput, poor coalescing, or L2/LDS pressure. |
AutoKernel Tier 2 focus: after launch geometry (Tier 1), maximize bytes moved usefully per cycle across HBM → L2 → L1/LDS → registers.
float4, uint4, Triton block_ptr with wide block_shape): fewer instructions, better line utilization when aligned.cp.async from global to shared; Hopper: TMA bulk tensor loads if using supported APIs/kernels.num_stages in Triton) to overlap load with math.+1 column) to skew addresses across banks.GROUP_SIZE_M) and problem splitting (Split-K, stream-K in Tier 4) change L2 reuse; pair Tier 2 vectorization with Tier 1 grouping.bank = (offset_in_bytes / 4) % 32 for classic layouts; 32 lanes hitting the same bank on the same cycle → serial replays.N_stride + 1 (FP32) is the minimal skew; larger pads sometimes help double-buffer pointer math stay conflict-free across stages.(row,col) to physical addresses so MMA loads spread across banks (see CUTLASS/Triton shared layouts).num_stages with buffer counts and verify bank conflict-free SMEM/LDS layouts (pad or swizzle).