L1 P27 · popcorn

correct 27_permutation_all_to_all.py

model

grok-4.3

turns / max

7 / 15

tool calls

tokens (total)

28575

wall clock

2652.0s

truncated

backend

cuda fp32

variant

popcorn

runtime

0.01 μs

ref runtime

0.03 μs

speedup

4.86x

started

2026-05-11T19:50:52Z

timingkernelbench.timing

kernel runtime

0.01 μs

ref runtime

0.03 μs

speedup

4.86x

kernel mean / median / std

0.01 / — / 1.61e-04 μs

ref mean / median / std

0.03 / — / 0.00 μ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

33.05 MB

ref peak

33.05 MB

kernel / ref

1.00x

kernel launches / fusiontorch.profiler

kernel launches

ref launches

fusion ratio (ref / kernel)

8.00x

total cuda time

23 μs

SOL (speed-of-light)nsight_hardware_counters

SOL score (0–1)

5.000e-04

dram util %

0.1%

compute util %

0.0%

bottleneck

unknown

dominant pipe

None

arithmetic intensity

—

ridge point

20.00

roofline / occupancynsight_hardware_counters

dram bandwidth

1.54 GB/s

dram util %

0.0%

fp32 tflops

—

fp32 util %

—

fp16 tflops

—

occupancy %

10.3%

bottleneck

unknown

dominant pipe %

—

arithmetic intensity

—

ridge point

20.00

L1 hit %

—

L2 hit %

99.3%

ld sectors / req

6.04

st sectors / req

—

regs / thread

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

3226.5%

short_scoreboard

1873.3%

reference PyTorch source (the prompt the agent sees)

"""
Fixed permutation along the feature axis: y[b, j] = x[b, perm[j]].
"""

import torch
import torch.nn as nn

from kernelbench.distributed_collectives import default_device

B = 10
N = 36

_g = torch.Generator()
_g.manual_seed(10)
PERM = torch.randperm(N, generator=_g)


class Model(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return x[:, PERM.to(device=x.device)]


def get_inputs():
    dev = default_device()
    return [torch.randn(popcorn_pri.jitter_int(B), N, device=dev)]


def get_init_inputs():
    return []

turn 0496.5s turn · 6.4s LLM · 490.1s tools · tokens in=1,937 · out=6 · total=2,362

reasoning

The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model does a fixed permutation along the feature axis: y[b, j] = x[b, perm[j]]. I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The original model permutes the columns of an input tensor x of shape (B, N) according to a fixed permutation tensor PERM, producing y = x[:, PERM].

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 1368.8s turn · 7.0s LLM · 361.8s tools · tokens in=2,595 · out=417 · total=3,096

reasoning

H100, high bandwidth. Sizes tiny: 10*36 floats = ~1.4KB, so memory bound trivially.

call → compile_kernel

callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 2370.3s turn · 7.6s LLM · 362.8s tools · tokens in=3,113 · out=418 · total=3,538

reasoning

Compiled ok. Now check correctness.

call → run_correctness

callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← run_correctness (fail)

failrun_correctness run_correctness FAILED: ? trials did not all match. runtime_error: expected scalar type Int but found Long

turn 3414.5s turn · 5.0s LLM · 409.5s tools · tokens in=3,567 · out=422 · total=4,035

reasoning

Error: perm is Long, but kernel expects int.

call → compile_kernel

callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.to(torch.int32).cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 4480.1s turn · 3.7s LLM · 476.4s tools · tokens in=4,052 · out=423 · total=4,478

reasoning

Now correctness.

call → run_correctness

callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.to(torch.int32).cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← run_correctness (ok)

okrun_correctness run_correctness PASSED: (3 / 3) 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 5365.8s turn · 4.6s LLM · 361.2s tools · tokens in=4,538 · out=422 · total=5,047

reasoning

Correct. Now, to check if faster, I need to submit or profile, but since submit ends, and to see speedup I should profile first?

call → profile_kernel

callprofile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.to(torch.int32).cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← profile_kernel (ok)

okprofile_kernel profile_kernel PASSED: profiling complete. === Kernel Profile Summary === --- Kernel Breakdown (by GPU time) --- #1: aten::empty 42.0μs (100%) calls=1 --- Timing --- GPU time: 3.39 μs --- Memory Bandwidth --- DRAM BW: 1.4/3350 GB/s (0.0%) read=1.4 write=0.0 GB/s Hit rates: L2=99.4% --- Compute Throughput --- Pipe util: FMA=1.0%, Tensor=0.1% --- Roofline Analysis --- Ridge point (FP32): 20.00 FLOPs/byte Classification: UNKNOWN --- Occupancy --- Achieved: 10.9% Eligible warps/cycle: 0.03 WARNING: <1 eligible warp/cycle = latency-bound, needs more parallelism Launch: regs/thread=16, smem/blk=1024B, block=256, grid=2 --- Memory Access Quality --- Global ld coalescing: 6.04 sectors/req (6.0x wasted BW) --- Warp Stall Reasons (% of issue cycles) --- long_scoreboard 2394.2% #################### short_scoreboard 2284.2% #################### --- Optimization Hints --- -> 2394% stalls: global memory latency (waiting for DRAM/L2). + sectors/req=6.0 -> uncoalesced global loads, restructure access pattern or use __ldg Try shared-memory tiling, prefetching, or reducing working set. -> Global load coalescing=6.0x (ideal=1.0) -> 6x wasted DRAM BW. Make per-warp accesses contiguous. -> Low occupancy (11%). Consider __launch_bounds__ or reducing register pressure. -> Low eligible warps + low BW% + low compute% = LATENCY-BOUND. Need more parallelism (larger grid, more ILP, prefetching).

turn 6156.0s turn · 8.2s LLM · 147.9s tools · tokens in=5,482 · out=422 · total=6,019 · FINAL

reasoning

The profile shows the kernel itself is not even timed? It shows aten::empty 42us, GPU time 3.39us? Weird, probably the kernel is tiny.

call → submit_kernel

callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N;\n if (idx < total) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[idx] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {\n int B = x.size(0);\n int N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int total = B * N;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);\n return y;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);\n\"\"\"\n\nmodule = load_inline(\n name=\"permute_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"permute_forward\"],\n verbose=False\n)\n\nPERM = None # will be set\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n global PERM\n if PERM is None:\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(36, generator=_g)\n self.perm = PERM.to(torch.int32).cuda()\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return module.permute_forward(x, self.perm)\n" }

result ← submit_kernel (ok)

oksubmit_kernel submit_kernel PASSED: (3 / 3) correctness trials all passed. Kernel runtime: 0.01 μs Runtime stats: mean=0.01μs median=?μs std=0.00μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 33.05 MB peak (ref: 33.05 MB, ratio: 1.0x) Kernel launches: 1 (ref: 8, fusion ratio: 8.0) SOL score: 0.0005

final submitted kernel (55 lines, 1,517 chars) — full source, untruncated

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

cuda_source = """
__global__ void permute_kernel(const float* __restrict__ x, float* __restrict__ y, const int* __restrict__ perm, int B, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * N;
    if (idx < total) {
        int b = idx / N;
        int j = idx % N;
        int p = perm[j];
        y[idx] = x[b * N + p];
    }
}

torch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm) {
    int B = x.size(0);
    int N = x.size(1);
    auto y = torch::empty({B, N}, x.options());
    int total = B * N;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), perm.data_ptr<int>(), B, N);
    return y;
}
"""

cpp_source = """
torch::Tensor permute_forward(torch::Tensor x, torch::Tensor perm);
"""

module = load_inline(
    name="permute_module",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["permute_forward"],
    verbose=False
)

PERM = None  # will be set

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        global PERM
        if PERM is None:
            _g = torch.Generator()
            _g.manual_seed(10)
            PERM = torch.randperm(36, generator=_g)
        self.perm = PERM.to(torch.int32).cuda()

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return module.permute_forward(x, self.perm)