L4 P1 · popcorn

error 1_meta-llama-Llama-3.2-1B-Instruct.py

model

Llama-4-Maverick-17B-128E-Instruct-FP8

turns / max

1 / 1

tool calls

tokens (total)

1431

wall clock

3.6s

truncated

backend

cuda fp32

variant

popcorn

runtime

—

ref runtime

0.00 μs

speedup

—

started

2026-05-11T08:41:15Z

reference PyTorch source (the prompt the agent sees)

import torch
from transformers import AutoConfig, AutoModelForCausalLM


class Model(torch.nn.Module):
    def __init__(self, model_name, config):
        super().__init__()
        self.model_name = model_name
        self.config = config
        self.model = AutoModelForCausalLM.from_pretrained(self.model_name, config=self.config)

    def forward(self, x):
        return self.model(x).logits


model_name = "meta-llama/Llama-3.2-1B-Instruct"
config = AutoConfig.from_pretrained(model_name)
vocab_size = config.vocab_size
sequence_length = 512
batch_size = 8


def get_inputs():
    p = popcorn_pri
    bs = p.jitter_int(batch_size)
    sl = p.jitter_int(sequence_length)
    inputs = torch.randint(0, vocab_size, (bs, sl))
    return [inputs]



def get_init_inputs():
    return [model_name, config]

turn 03.6s turn · 3.5s LLM · 0.1s tools · tokens in=961 · out=470 · total=1,431

call → submit_kernel

callsubmit_kernel { "kernel_code": "import torch\nfrom torch.utils.cpp_extension import load_inline\n\n# Define the CUDA kernel\ncuda_source = '''\n#include <cuda_runtime.h>\n#include <cuda_fp16.h>\n\n__global__ void custom_kernel(float* output, int* input, int vocab_size, int sequence_length, int batch_size) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx < batch_size * sequence_length) {\n int input_val = input[idx];\n if (input_val >= 0 && input_val < vocab_size) {\n output[idx] = (float)input_val;\n } else {\n output[idx] = 0.0f;\n }\n }\n}\n\nvoid wrapper(float* output, int* input, int vocab_size, int sequence_length, int batch_size) {\n int blockSize = 256;\n int numBlocks = (batch_size * sequence_length + blockSize - 1) / blockSize;\n custom_kernel<<<numBlocks, blockSize>>>(output, input, vocab_size, sequence_length, batch_size);\n cudaDeviceSynchronize();\n}\n'''\n\n# Compile the CUDA kernel\ncustom_kernel_module = load_inline(\n name='custom_kernel_module',\n cuda_sources=[cuda_source],\n functions=['wrapper']\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self, model_name, config):\n super().__init__()\n self.model_name = model_name\n self.config = config\n self.vocab_size = config.vocab_size\n\n def forward(self, x):\n batch_size, sequence_length = x.shape\n output = torch.empty(batch_size, sequence_length, dtype=torch.float32, device=x.device)\n custom_kernel_module.wrapper(\n output.data_ptr(),\n x.data_ptr(),\n self.vocab_size,\n sequence_length,\n batch_size\n )\n return output\n" }

result ← submit_kernel (fail)

failsubmit_kernel submit_kernel FAILED: TypeError: cannot unpack non-iterable NoneType object

final submitted kernel (54 lines, 1,656 chars) — full source, untruncated

import torch
from torch.utils.cpp_extension import load_inline

# Define the CUDA kernel
cuda_source = '''
#include <cuda_runtime.h>
#include <cuda_fp16.h>

__global__ void custom_kernel(float* output, int* input, int vocab_size, int sequence_length, int batch_size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < batch_size * sequence_length) {
        int input_val = input[idx];
        if (input_val >= 0 && input_val < vocab_size) {
            output[idx] = (float)input_val;
        } else {
            output[idx] = 0.0f;
        }
    }
}

void wrapper(float* output, int* input, int vocab_size, int sequence_length, int batch_size) {
    int blockSize = 256;
    int numBlocks = (batch_size * sequence_length + blockSize - 1) / blockSize;
    custom_kernel<<<numBlocks, blockSize>>>(output, input, vocab_size, sequence_length, batch_size);
    cudaDeviceSynchronize();
}
'''

# Compile the CUDA kernel
custom_kernel_module = load_inline(
    name='custom_kernel_module',
    cuda_sources=[cuda_source],
    functions=['wrapper']
)

class ModelNew(torch.nn.Module):
    def __init__(self, model_name, config):
        super().__init__()
        self.model_name = model_name
        self.config = config
        self.vocab_size = config.vocab_size

    def forward(self, x):
        batch_size, sequence_length = x.shape
        output = torch.empty(batch_size, sequence_length, dtype=torch.float32, device=x.device)
        custom_kernel_module.wrapper(
            output.data_ptr(),
            x.data_ptr(),
            self.vocab_size,
            sequence_length,
            batch_size
        )
        return output