Kimi-K2.6 · original · l12_default_gpt
correct 27_SELU_.py
model
Kimi-K2.6
turns / max
5 / 10
tool calls
5
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
original
runtime
4.34 μs
ref runtime
4.30 μs
speedup
0.99x
started
2026-05-02T02:14:35Z
timingkernelbench.timing
kernel runtime
4.34 μs
ref runtime
4.30 μs
speedup
0.99x
kernel mean / median / std
4.34 / — / 0.08 μs
ref mean / median / std
4.30 / — / 0.09 μs
numerical precisioneval_kernel_against_ref
max abs error
0.00
mean abs error
0.00
max rel error
0.00
mean rel error
0.00
gpu memoryextended_metrics.measure_memory
kernel peak
12288.00 MB
ref peak
12288.00 MB
kernel / ref
1.00x
kernel launches / fusiontorch.profiler
kernel launches
2
ref launches
2
fusion ratio (ref / kernel)
1.00x
total cuda time
94 μs
energyNVML (pynvml)
kernel mJ / run
2417.60
ref mJ / run
1953.18
ratio (ref / kernel)
0.81x
avg power
559.07 W
ref avg power
449.10 W
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.904
dram util %
90.4%
compute util %
5.3%
bottleneck
memory-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
3029.81 GB/s
dram util %
90.4%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
78.7%
bottleneck
memory-bound
dominant pipe %
5.4%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
50.5%
ld sectors / req
16.00
st sectors / req
—
regs / thread
21
smem / block
1024 B
block size
256
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
8507.6%
short_scoreboard
109.4%
mio_throttle
3.7%
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs a SELU activation.
    """
    def __init__(self):
        super(Model, self).__init__()
    
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        """
        Applies SELU activation to the input tensor.

        Args:
            x (torch.Tensor): Input tensor of any shape.

        Returns:
            torch.Tensor: Output tensor with SELU applied, same shape as input.
        """
        return torch.selu(x)
    
batch_size = 4096
dim = 393216

def get_inputs():
    x = torch.rand(batch_size, dim)
    return [x]

def get_init_inputs():
    return []  # No special initialization inputs needed
turn 05.2s turn · 5.2s LLM · 0.0s tools
reasoning
Let's analyze the problem: 1. We need to replace `torch.selu(x)` with a custom CUDA kernel. 2. The input shape is `[batch_size, dim] = [4096, 393216]`. 3. This is a large tensor with ~1.6 billion elements (4096 * 393216 = 1,610,612,736). 4. Each element is independent, so this is an element-wise kernel. 5. For large element-wise operations, performance is memory bandwidth bound. 6. We need to write an efficient fused CUDA kernel that loads inputs, applies SELU, and stores outputs. 7. SELU formula: `scale * (max(0, x) + min(0, alpha * (exp(x) - 1)))` where `scale ≈ 1.0507` and `alpha ≈ 1.6733`. Let me check the exact PyTorch SELU parameters: - alpha = 1.6732632423543772848170429916717 - scale = 1.0507010221481323242187531591474 The formula is: ``` selu(x) = scale * x if x > 0 selu(x) = scale * alpha * (exp(x) - 1) if x <= 0 ``` For a kernel of 1.6B elements, we need to maximize memory throughput. Standard approaches: - Use vectorized loads/stores (float4) - Use grid-stride loops - Ensure coalesced memory access Let me get the GPU specs first to understand memory bandwidth and optimal block size.
call → get_gpu_specs
callget_gpu_specs {}
result ← get_gpu_specs (ok)
okget_gpu_specs GPU specs for NVIDIA H100 80GB HBM3: total memory (runtime): 79.1 GB spec entry: H100 GPU Architecture: Hopper GPU Memory: 80GB Memory Bandwidth: 3.35 TB/s FP64 TFLOPS: 34 FP64 Tensor Core TFLOPS: 67 FP32 TFLOPS: 67 TF32 Tensor Core TFLOPS: 989 with sparsity BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity FP16 Tensor Core TFLOPS: 1979 with sparsity FP8 Tensor Core TFLOPS: 3958 with sparsity INT8 Tensor Core TOPS: 3958 with sparsity Register File Size: 64K 32-bit registers per SM Maximum number of registers per thread: 255 Maximum number of thread blocks per SM: 32 Shared memory capacity per SM: 228 KB Maximum shared memory per thread block: 227 KB
turn 1142.7s turn · 142.7s LLM · 0.0s tools
reasoning
The task is to replace `torch.selu` with a custom CUDA kernel for H100. Let me analyze what's needed: 1. Input: [4096, 393216] = 1,610,612,736 elements (~1.6 billion) 2. `torch.selu(x)` is a simple element-wise operation: `scale * max(0, x) + scale * alpha * (exp(min(0, x)) - 1)` where scale ≈ 1.0507, alpha ≈ 1.67326 For H100 with 3.35 TB/s bandwidth, this is clearly memory-bandwidth bound. The input is fp32 (4 bytes), output is fp32 (4 bytes). Total memory traffic is ~12.88 GB. At 3.35 TB/s, theoretical minimum time is ~3.84 ms. To maximize performance, I need to: 1. Use a simple CUDA kernel with vectorized loads/stores (e.g., float4) 2. Use enough threads per block (256 or 512) 3. Ensure grid size is large enough to saturate the GPU Let me write a kernel: - Process 4 elements per thread using float4 - Grid-stride loop For SELU constants: - scale = 1.0507009873554804934193349852946 - alpha = 1.6732632423543772848180429911717 Formula: `scale * x` if x > 0 else `scale * alpha * (exp(x) - 1)` Let me write the kernel code.
call → compile_kernel
callcompile_kernel { "kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2271.0s turn · 271.0s LLM · 0.0s tools
reasoning
Good, it compiled. Now let's check correctness.
call → run_correctness
callrun_correctness { "kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 3101.1s turn · 101.1s LLM · 0.0s tools
reasoning
The kernel compiles and passes correctness. Let me now run a static analysis before submitting.
call → static_check
callstatic_check { "kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n" }
result ← static_check (ok)
okstatic_check static_check PASSED: no violations or warnings detected.
turn 429.8s turn · 29.8s LLM · 0.0s tools · FINAL
reasoning
The kernel passes correctness and static checks. Let me now submit it. But before I submit, let me review once more: 1. The kernel uses float4 vectorized loads/stores with proper alignment considerations. 2. The SELU formula is correctly implemented with the right constants. 3. The kernel handles the tail case where n is not a multiple of 4. 4. The block size is 256 which is good. 5. `-O3` and `--use_fast_math` flags are used for optimization. However, I should double-check: the input `x = torch.rand(batch_size, dim)` will be in the range [0, 1), so all values are positive. This means the SELU will always take the positive branch. Does my kernel handle this correctly? Yes, because `x > 0` evaluates to true, and we return `scale * x`. One concern: `float4` requires 16-byte alignment. For tensors allocated by PyTorch, this should generally be satisfied since CUDA memory allocations are typically aligned to at least 256 bytes. And `batch_size * dim = 4096 * 393216 = 1,610,612,736`, which is divisible by 4 (and by 16), so the total number of elements is a multiple of 4. Actually, wait. The starting address of the tensor data should be aligned to at least 16 bytes for float4. PyTorch CUDA tensors are allocated with 256-byte alignment typically, so this should be fine. Let me submit the kernel.
call → submit_kernel
callsubmit_kernel { "kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 4.34 μs Runtime stats: mean=4.34μs median=?μs std=0.08μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 12288.0 MB peak (ref: 12288.0 MB, ratio: 1.0x) Kernel launches: 2 (ref: 2, fusion ratio: 1.0) Energy: 2417.6 mJ/run (ref: 1953.18 mJ, ratio: 0.8079x) SOL score: 0.9044
final submitted kernel (70 lines, 2,251 chars) — full source, untruncated
import torch
from torch.utils.cpp_extension import load_inline

cuda_source = '''
#include <torch/extension.h>
#include <cuda_runtime.h>

constexpr float SELU_SCALE = 1.0507009873554804934193349852946f;
constexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;

__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {
    constexpr int VEC_SIZE = 4;
    int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;
    
    if (idx + VEC_SIZE <= n) {
        float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];
        float4 out;
        
        out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);
        out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);
        out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);
        out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);
        
        reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;
    } else {
        for (int64_t i = idx; i < n; ++i) {
            float x = input[i];
            output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);
        }
    }
}

torch::Tensor selu_forward(torch::Tensor input) {
    auto output = torch::empty_like(input);
    int64_t n = input.numel();
    
    if (n == 0) return output;
    
    constexpr int VEC_SIZE = 4;
    constexpr int BLOCK_SIZE = 256;
    int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;
    int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;
    
    selu_kernel<<<num_blocks, BLOCK_SIZE>>>(
        input.data_ptr<float>(),
        output.data_ptr<float>(),
        n
    );
    
    return output;
}
'''

selu_module = load_inline(
    name='selu_cuda',
    cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',
    cuda_sources=cuda_source,
    functions=['selu_forward'],
    extra_cuda_cflags=['-O3', '--use_fast_math'],
    verbose=False
)

class ModelNew(torch.nn.Module):
    def __init__(self):
        super(ModelNew, self).__init__()
    
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return selu_module.selu_forward(x)