grok-4.3 · popcorn · pop_l123_all_gpt
correct 16_DeepSeekMoECombineScatter.py
model
grok-4.3
turns / max
5 / 15
tool calls
5
tokens (total)
21033
wall clock
976.8s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.07 μs
ref runtime
97.60 μs
speedup
1469.88x
started
2026-05-11T18:33:55Z
timingkernelbench.timing
kernel runtime
0.07 μs
ref runtime
97.60 μs
speedup
1469.88x
kernel mean / median / std
0.07 / — / 0.05 μs
ref mean / median / std
97.60 / — / 0.54 μ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
35.30 MB
ref peak
35.27 MB
kernel / ref
1.00x
kernel launches / fusiontorch.profiler
kernel launches
8
ref launches
46084
fusion ratio (ref / kernel)
5760.50x
total cuda time
315 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.022
dram util %
2.2%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
72.42 GB/s
dram util %
2.2%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
27.1%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
285.0%
ld sectors / req
35.60
st sectors / req
—
regs / thread
80
smem / block
3072 B
block size
512
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
10140.8%
short_scoreboard
800.9%
flagssubmit_kernel guard
excessive speedup
flagged
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Weighted expert-output combine back into token-major order.

    Expert outputs are scattered back to token positions and scaled by the
    routing weight associated with each expert contribution.
    """

    def __init__(self):
        super().__init__()

    def forward(
        self,
        expert_hidden: torch.Tensor,
        token_idx: torch.Tensor,
        gates: torch.Tensor,
        num_tokens: int,
    ) -> torch.Tensor:
        out = torch.zeros(num_tokens, expert_hidden.shape[1], dtype=expert_hidden.dtype, device=expert_hidden.device)
        for row in range(expert_hidden.shape[0]):
            token = int(token_idx[row].item())
            out[token] += gates[row] * expert_hidden[row]
        return out


num_tokens = 1536
hidden_dim = 128
fanout = 2


def get_inputs():
    return popcorn_pri.moe_combine_inputs(num_tokens, hidden_dim, fanout)




def get_init_inputs():
    return []
turn 0375.6s turn · 19.6s LLM · 356.0s tools · tokens in=2,046 · out=475 · total=3,058
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model combines expert outputs back into token-major order, weighted by routing gates. The task requires replacing the forward computation of a PyTorch model with a custom CUDA kernel in a ModelNew class. The goal is to efficiently recombine expert outputs into token-major order, weighted by routing gates. The reference implementation uses a loop to scatter-add scaled expert_hiddens back into an output tensor.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,\n const int* __restrict__ token_idx,\n const float* __restrict__ gates,\n float* __restrict__ out,\n int N, int D) {\n int row = blockIdx.x * blockDim.x + threadIdx.x;\n if (row < N) {\n int token = token_idx[row];\n float gate = gates[row];\n const float* src = expert_hidden + row * D;\n float* dst = out + token * D;\n for (int d = 0; d < D; d++) {\n atomicAdd(&dst[d], gate * src[d]);\n }\n }\n}\n\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens) {\n int N = expert_hidden.size(0);\n int D = expert_hidden.size(1);\n auto out = torch::zeros({num_tokens, D}, expert_hidden.options());\n const int block_size = 256;\n const int num_blocks = (N + block_size - 1) / block_size;\n scatter_add_kernel<<<num_blocks, block_size>>>(\n expert_hidden.data_ptr<float>(),\n token_idx.data_ptr<int>(),\n gates.data_ptr<float>(),\n out.data_ptr<float>(),\n N, D\n );\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens);\n'''\n\nmodule = load_inline(\n name='scatter_add',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scatter_add_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, expert_hidden, token_idx, gates, num_tokens):\n return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 1107.9s turn · 6.2s LLM · 101.7s tools · tokens in=3,075 · out=476 · total=3,558
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.h>\n#include <cuda_runtime.h>\n\n__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,\n const int* __restrict__ token_idx,\n const float* __restrict__ gates,\n float* __restrict__ out,\n int N, int D) {\n int row = blockIdx.x * blockDim.x + threadIdx.x;\n if (row < N) {\n int token = token_idx[row];\n float gate = gates[row];\n const float* src = expert_hidden + row * D;\n float* dst = out + token * D;\n for (int d = 0; d < D; d++) {\n atomicAdd(&dst[d], gate * src[d]);\n }\n }\n}\n\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens) {\n int N = expert_hidden.size(0);\n int D = expert_hidden.size(1);\n auto out = torch::zeros({num_tokens, D}, expert_hidden.options());\n const int block_size = 256;\n const int num_blocks = (N + block_size - 1) / block_size;\n scatter_add_kernel<<<num_blocks, block_size>>>(\n expert_hidden.data_ptr<float>(),\n token_idx.data_ptr<int>(),\n gates.data_ptr<float>(),\n out.data_ptr<float>(),\n N, D\n );\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens);\n'''\n\nmodule = load_inline(\n name='scatter_add',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scatter_add_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, expert_hidden, token_idx, gates, num_tokens):\n return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)\n" }
result ← run_correctness (fail)
failrun_correctness run_correctness FAILED: ? trials did not all match. runtime_error: expected scalar type Int but found Float
turn 2150.3s turn · 8.9s LLM · 141.5s tools · tokens in=3,587 · out=492 · total=4,275
reasoning
Error: token_idx is int32 but probably needs to be int64 or something? In PyTorch, indices are often long.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,\n const int64_t* __restrict__ token_idx,\n const float* __restrict__ gates,\n float* __restrict__ out,\n int N, int D) {\n int row = blockIdx.x * blockDim.x + threadIdx.x;\n if (row < N) {\n int64_t token = token_idx[row];\n float gate = gates[row];\n const float* src = expert_hidden + row * D;\n float* dst = out + token * D;\n for (int d = 0; d < D; d++) {\n atomicAdd(&dst[d], gate * src[d]);\n }\n }\n}\n\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens) {\n int N = expert_hidden.size(0);\n int D = expert_hidden.size(1);\n auto out = torch::zeros({num_tokens, D}, expert_hidden.options());\n const int block_size = 256;\n const int num_blocks = (N + block_size - 1) / block_size;\n scatter_add_kernel<<<num_blocks, block_size>>>(\n expert_hidden.data_ptr<float>(),\n token_idx.data_ptr<int64_t>(),\n gates.data_ptr<float>(),\n out.data_ptr<float>(),\n N, D\n );\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens);\n'''\n\nmodule = load_inline(\n name='scatter_add',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scatter_add_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, expert_hidden, token_idx, gates, num_tokens):\n token_idx = token_idx.to(torch.int64)\n return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 3166.0s turn · 4.8s LLM · 161.2s tools · tokens in=4,292 · out=493 · total=4,789
reasoning
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.h>\n#include <cuda_runtime.h>\n\n__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,\n const int64_t* __restrict__ token_idx,\n const float* __restrict__ gates,\n float* __restrict__ out,\n int N, int D) {\n int row = blockIdx.x * blockDim.x + threadIdx.x;\n if (row < N) {\n int64_t token = token_idx[row];\n float gate = gates[row];\n const float* src = expert_hidden + row * D;\n float* dst = out + token * D;\n for (int d = 0; d < D; d++) {\n atomicAdd(&dst[d], gate * src[d]);\n }\n }\n}\n\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens) {\n int N = expert_hidden.size(0);\n int D = expert_hidden.size(1);\n auto out = torch::zeros({num_tokens, D}, expert_hidden.options());\n const int block_size = 256;\n const int num_blocks = (N + block_size - 1) / block_size;\n scatter_add_kernel<<<num_blocks, block_size>>>(\n expert_hidden.data_ptr<float>(),\n token_idx.data_ptr<int64_t>(),\n gates.data_ptr<float>(),\n out.data_ptr<float>(),\n N, D\n );\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens);\n'''\n\nmodule = load_inline(\n name='scatter_add',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scatter_add_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, expert_hidden, token_idx, gates, num_tokens):\n token_idx = token_idx.to(torch.int64)\n return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)\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 4177.0s turn · 5.9s LLM · 171.1s tools · tokens in=4,849 · out=492 · total=5,353 · FINAL
reasoning
Correct! Now, to submit, since it's faster presumably.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,\n const int64_t* __restrict__ token_idx,\n const float* __restrict__ gates,\n float* __restrict__ out,\n int N, int D) {\n int row = blockIdx.x * blockDim.x + threadIdx.x;\n if (row < N) {\n int64_t token = token_idx[row];\n float gate = gates[row];\n const float* src = expert_hidden + row * D;\n float* dst = out + token * D;\n for (int d = 0; d < D; d++) {\n atomicAdd(&dst[d], gate * src[d]);\n }\n }\n}\n\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens) {\n int N = expert_hidden.size(0);\n int D = expert_hidden.size(1);\n auto out = torch::zeros({num_tokens, D}, expert_hidden.options());\n const int block_size = 256;\n const int num_blocks = (N + block_size - 1) / block_size;\n scatter_add_kernel<<<num_blocks, block_size>>>(\n expert_hidden.data_ptr<float>(),\n token_idx.data_ptr<int64_t>(),\n gates.data_ptr<float>(),\n out.data_ptr<float>(),\n N, D\n );\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,\n torch::Tensor token_idx,\n torch::Tensor gates,\n int num_tokens);\n'''\n\nmodule = load_inline(\n name='scatter_add',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scatter_add_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, expert_hidden, token_idx, gates, num_tokens):\n token_idx = token_idx.to(torch.int64)\n return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (3 / 3) correctness trials all passed. Kernel runtime: 0.07 μs Runtime stats: mean=0.07μs median=?μs std=0.05μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 35.3 MB peak (ref: 35.27 MB, ratio: 1.0007x) Kernel launches: 8 (ref: 46084, fusion ratio: 5760.5) SOL score: 0.0216 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (70 lines, 2,255 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.h>
#include <cuda_runtime.h>

__global__ void scatter_add_kernel(const float* __restrict__ expert_hidden,
                                    const int64_t* __restrict__ token_idx,
                                    const float* __restrict__ gates,
                                    float* __restrict__ out,
                                    int N, int D) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N) {
        int64_t token = token_idx[row];
        float gate = gates[row];
        const float* src = expert_hidden + row * D;
        float* dst = out + token * D;
        for (int d = 0; d < D; d++) {
            atomicAdd(&dst[d], gate * src[d]);
        }
    }
}

torch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,
                               torch::Tensor token_idx,
                               torch::Tensor gates,
                               int num_tokens) {
    int N = expert_hidden.size(0);
    int D = expert_hidden.size(1);
    auto out = torch::zeros({num_tokens, D}, expert_hidden.options());
    const int block_size = 256;
    const int num_blocks = (N + block_size - 1) / block_size;
    scatter_add_kernel<<<num_blocks, block_size>>>(
        expert_hidden.data_ptr<float>(),
        token_idx.data_ptr<int64_t>(),
        gates.data_ptr<float>(),
        out.data_ptr<float>(),
        N, D
    );
    return out;
}
'''

cpp_source = '''
torch::Tensor scatter_add_cuda(torch::Tensor expert_hidden,
                               torch::Tensor token_idx,
                               torch::Tensor gates,
                               int num_tokens);
'''

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

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

    def forward(self, expert_hidden, token_idx, gates, num_tokens):
        token_idx = token_idx.to(torch.int64)
        return self.module.scatter_add_cuda(expert_hidden, token_idx, gates, num_tokens)