L1 P27 · popcorn

correct 27_permutation_all_to_all.py

model

grok-4.3

turns / max

9 / 12

tool calls

tokens (total)

34544

wall clock

1146.3s

truncated

backend

cuda fp32

variant

popcorn

runtime

0.03 μs

ref runtime

0.04 μs

speedup

1.23x

started

2026-05-11T12:32:47Z

timingkernelbench.timing

kernel runtime

0.03 μs

ref runtime

0.04 μs

speedup

1.23x

kernel mean / median / std

0.03 / — / 0.00 μs

ref mean / median / std

0.04 / — / 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.00 MB

ref peak

33.00 MB

kernel / ref

1.00x

kernel launches / fusiontorch.profiler

kernel launches

ref launches

fusion ratio (ref / kernel)

1.60x

total cuda time

217 μs

SOL (speed-of-light)nsight_hardware_counters

SOL score (0–1)

4.000e-04

dram util %

0.0%

compute util %

0.0%

bottleneck

unknown

dominant pipe

None

arithmetic intensity

—

ridge point

20.00

roofline / occupancynsight_hardware_counters

dram bandwidth

1.45 GB/s

dram util %

0.0%

fp32 tflops

—

fp32 util %

—

fp16 tflops

—

occupancy %

10.6%

bottleneck

unknown

dominant pipe %

—

arithmetic intensity

—

ridge point

20.00

L1 hit %

—

L2 hit %

99.3%

ld sectors / req

7.50

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

2856.6%

short_scoreboard

1703.8%

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 076.1s turn · 2.7s LLM · 73.4s tools · tokens in=1,406 · out=6 · total=1,514

reasoning

The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs a fixed permutation along the feature axis: y[b, j] = x[b, perm[j]].

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 1112.4s turn · 11.5s LLM · 100.9s tools · tokens in=1,747 · out=438 · total=2,468

reasoning

The operation is a permutation of columns in a (B, N) matrix, with B=10, N=36, very small. But we need to make it faster than PyTorch.

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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n from kernelbench.distributed_collectives import default_device\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g, device=x.device)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 2272.3s turn · 3.9s LLM · 268.4s tools · tokens in=2,485 · out=439 · total=2,931

reasoning

Compiled successfully. 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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n from kernelbench.distributed_collectives import default_device\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g, device=x.device)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← run_correctness (fail)

failrun_correctness run_correctness FAILED: ? trials did not all match. runtime_error: Expected a 'cuda' device type for generator but found 'cpu'

turn 3129.6s turn · 3.9s LLM · 125.8s tools · tokens in=2,967 · out=422 · total=3,448

reasoning

Error because generator device. Need to fix the perm creation.

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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 4103.7s turn · 2.7s LLM · 100.9s tools · tokens in=3,465 · out=423 · total=3,892

reasoning

Now run 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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\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 5136.9s turn · 2.7s LLM · 134.3s tools · tokens in=3,921 · out=428 · total=4,377

reasoning

Perm is long, need int64 or cast to 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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int64_t* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int64_t p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int64_t>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← compile_kernel (ok)

okcompile_kernel compile_kernel PASSED: kernel compiled without errors.

turn 6108.0s turn · 2.9s LLM · 105.1s tools · tokens in=4,394 · out=429 · total=4,826

reasoning

Now test.

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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int64_t* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int64_t p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int64_t>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\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 74.8s turn · 4.8s LLM · 0.0s tools · tokens in=4,886 · out=428 · total=5,319

reasoning

Correct! Now submit.

call → static_check

callstatic_check { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int64_t* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int64_t p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int64_t>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← static_check (ok)

okstatic_check static_check PASSED: no violations or warnings detected.

turn 8202.4s turn · 5.8s LLM · 196.6s tools · tokens in=5,337 · out=428 · total=5,769 · FINAL

reasoning

Good, submit.

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#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void permute_kernel(const float* __restrict__ x, const int64_t* __restrict__ perm, float* __restrict__ y, int B, int N) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < B * N) {\n int b = idx / N;\n int j = idx % N;\n int64_t p = perm[j];\n y[b * N + j] = x[b * N + p];\n }\n}\n\ntorch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm) {\n auto B = x.size(0);\n auto N = x.size(1);\n auto y = torch::empty({B, N}, x.options());\n int threads = 256;\n int blocks = (B * N + threads - 1) / threads;\n permute_kernel<<<blocks, threads>>>(x.data_ptr<float>(), perm.data_ptr<int64_t>(), y.data_ptr<float>(), B, N);\n return y;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor permute_cuda(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_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.perm = None\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n if self.perm is None:\n import torch\n _g = torch.Generator()\n _g.manual_seed(10)\n PERM = torch.randperm(x.size(1), generator=_g)\n self.perm = PERM\n return module.permute_cuda(x, self.perm.to(x.device))\n" }

result ← submit_kernel (ok)

oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.03 μs Runtime stats: mean=0.03μ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.0 MB peak (ref: 33.0 MB, ratio: 1.0x) Kernel launches: 5 (ref: 8, fusion ratio: 1.6) SOL score: 0.0004

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

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

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

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

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

cpp_source = '''
torch::Tensor permute_cuda(torch::Tensor x, torch::Tensor perm);
'''

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

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self.perm = None

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        if self.perm is None:
            import torch
            _g = torch.Generator()
            _g.manual_seed(10)
            PERM = torch.randperm(x.size(1), generator=_g)
            self.perm = PERM
        return module.permute_cuda(x, self.perm.to(x.device))