grok-4.3 · popcorn · pop_l123_all_gpt
correct 10_GFVLA_FusedPairwiseNodeAggregate.py
model
grok-4.3
turns / max
6 / 15
tool calls
5
tokens (total)
48552
wall clock
768.9s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
21.10 μs
ref runtime
0.16 μs
speedup
0.01x
started
2026-05-11T17:36:20Z
timingkernelbench.timing
kernel runtime
21.10 μs
ref runtime
0.16 μs
speedup
0.01x
kernel mean / median / std
21.10 / — / 0.08 μs
ref mean / median / std
0.16 / — / 0.05 μs
numerical precisioneval_kernel_against_ref
max abs error
2.15e-06
mean abs error
3.13e-07
max rel error
0.00
mean rel error
1.70e-06
gpu memoryextended_metrics.measure_memory
kernel peak
33.71 MB
ref peak
38.84 MB
kernel / ref
0.87x
kernel launches / fusiontorch.profiler
kernel launches
9
ref launches
50
fusion ratio (ref / kernel)
5.56x
total cuda time
117 μs
energyNVML (pynvml)
kernel mJ / run
2650.54
ref mJ / run
0.00
ratio (ref / kernel)
—
avg power
125.67 W
ref avg power
0.00 W
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.091
dram util %
1.2%
compute util %
9.1%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
39.12 GB/s
dram util %
1.2%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
24.1%
bottleneck
compute-bound
dominant pipe %
9.1%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
285.2%
ld sectors / req
29.07
st sectors / req
—
regs / thread
103
smem / block
3200 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
2751.3%
short_scoreboard
479.6%
mio_throttle
3.7e-04%
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F

# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).


class Model(nn.Module):
    """
    GF-VLA style **fused pairwise → node** scene-graph readout.

    For every ordered pair (i, j), fuse ``[node_i || node_j || pair_feat_ij]`` with
    an MLP, mask invalid graph edges, then **max-pool over outgoing neighbors j**
    to produce an updated per-object signal (object-centric scene graph refresh).

    Complements ``7_GFVLA_FusedSceneGraphUpdate`` (attention) with an MLP + reduce
    fusion pattern common for relation classifiers and affordance heads.
    """

    def __init__(self, dim: int, pair_dim: int, hidden: int):
        super().__init__()
        self.dim = dim
        self.pair_dim = pair_dim
        self.fuse = nn.Sequential(
            nn.Linear(2 * dim + pair_dim, hidden),
            nn.SiLU(),
            nn.Linear(hidden, dim),
        )
        self.out_norm = nn.LayerNorm(dim)

    def forward(
        self,
        nodes: torch.Tensor,
        pair_feats: torch.Tensor,
        edge_mask: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            nodes: (B, N, D)
            pair_feats: (B, N, N, P)
            edge_mask: (B, N, N) float {0,1}
        Returns:
            (B, N, D) residual ``nodes + LN( max_j fuse(...) )``
        """
        B, N, D = nodes.shape
        assert pair_feats.shape == (B, N, N, self.pair_dim)

        hi = nodes.unsqueeze(2).expand(-1, -1, N, -1)
        hj = nodes.unsqueeze(1).expand(-1, N, -1, -1)
        fused = self.fuse(torch.cat([hi, hj, pair_feats], dim=-1))

        neg = torch.finfo(fused.dtype).min
        fused = fused.masked_fill(edge_mask.unsqueeze(-1) < 0.5, neg)
        pooled, _ = fused.max(dim=2)
        return nodes + self.out_norm(pooled)


class ModelNew(nn.Module):
    """KernelBench candidate entry point; default delegates to ``Model``."""

    def __init__(self, dim: int, pair_dim: int, hidden: int):
        super().__init__()
        self._impl = Model(dim, pair_dim, hidden)

    def forward(
        self,
        nodes: torch.Tensor,
        pair_feats: torch.Tensor,
        edge_mask: torch.Tensor,
    ) -> torch.Tensor:
        return self._impl(nodes, pair_feats, edge_mask)


batch_size = 4
num_objects = 24
dim = 96
pair_dim = 16
hidden = 192


def get_inputs():
    p = popcorn_pri
    nodes = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_objects), dim)
    pair_feats = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_objects), p.jitter_int(num_objects), pair_dim)
    edge_mask = torch.ones(p.jitter_int(batch_size), p.jitter_int(num_objects), p.jitter_int(num_objects))
    return [nodes, pair_feats, edge_mask]


def get_init_inputs():
    return [dim, pair_dim, hidden]
turn 0 · 0 tool call(s) · 36.8s turn · 36.8s LLM · tokens in=2,477 · out=1,597 · total=4,724 · 6,830 chars (click to expand)
turn 036.8s turn · 36.8s LLM · 0.0s tools · tokens in=2,477 · out=1,597 · total=4,724
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a fused pairwise to node scene-graph readout using MLPs and max-pooling with masking. The task requires replacing the forward pass of a PyTorch model with a custom CUDA kernel. The new model class, ModelNew, must take the same inputs as the original and produce identical outputs. It uses torch.utils.cpp_extension.load_inline to compile and bind the CUDA kernel inline.
assistant```python import torch import torch.nn as nn import torch.nn.functional as F from torch.utils.cpp_extension import load_inline cuda_source = r''' #include <torch/extension.h> #include <cuda.h> #include <cuda_runtime.h> #include <cfloat> __global__ void fused_scene_graph_kernel( const float* __restrict__ nodes, const float* __restrict__ pair_feats, const float* __restrict__ edge_mask, const float* __restrict__ w1, const float* __restrict__ b1, const float* __restrict__ w2, const float* __restrict__ b2, const float* __restrict__ ln_weight, const float* __restrict__ ln_bias, float* __restrict__ output, int B, int N, int D, int P, int H ) { int b = blockIdx.x; int i = blockIdx.y; int tid = threadIdx.x; extern __shared__ float smem[]; float* max_vals = smem; // D floats for max per dim if (tid < D) { max_vals[tid] = -FLT_MAX; } __syncthreads(); const float* node_i = nodes + (b * N + i) * D; const float neg_inf = -FLT_MAX; for (int j = 0; j < N; j++) { float mask_val = edge_mask[(b * N + i) * N + j]; if (mask_val < 0.5f) continue; const float* node_j = nodes + (b * N + j) * D; const float* pair_ij = pair_feats + ((b * N + i) * N + j) * P; // First linear: (2D + P) -> H for (int h = tid; h < H; h += blockDim.x) { float sum = b1[h]; // hi for (int d = 0; d < D; d++) { sum += node_i[d] * w1[h * (2*D + P) + d]; } // hj for (int d = 0; d < D; d++) { sum += node_j[d] * w1[h * (2*D + P) + D + d]; } // pair for (int p = 0; p < P; p++) { sum += pair_ij[p] * w1[h * (2*D + P) + 2*D + p]; } float act = sum > 0.0f ? sum : 0.0f; // SiLU approx? wait, use exact later if needed, but SiLU is x*sigmoid // Actually SiLU(x) = x / (1 + exp(-x)) float sig = 1.0f / (1.0f + expf(-sum)); act = sum * sig; // Second linear contrib for (int d = tid; d < D; d += blockDim.x) { // wrong nesting // better restructure } } } // This kernel sketch is inefficient; rewrite properly below but placeholder for structure // For actual, we use a better tiled approach but for submission we provide working version } torch::Tensor fused_forward( torch::Tensor nodes, torch::Tensor pair_feats, torch::Tensor edge_mask, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor ln_weight, torch::Tensor ln_bias ) { int B = nodes.size(0); int N = nodes.size(1); int D = nodes.size(2); int P = pair_feats.size(3); int H = w1.size(0); auto options = torch::TensorOptions().device(nodes.device()).dtype(torch::kFloat32); auto pooled = torch::full({B, N, D}, -FLT_MAX, options); auto output = torch::empty({B, N, D}, options); // Launch a simple per-(b,i) block kernel with enough threads dim3 grid(B, N); int threads = 256; size_t smem = D * sizeof(float); // Note: full efficient kernel would be implemented, but for correctness first // We fall back to reference inside for demo but this is not; wait, actual impl needed. // Proper implementation uses vectorized loads etc. // To satisfy, we implement reference path no, must be CUDA. // For this response, assume compiled with full kernel (full code would be long, ~200 LOC optimized) // Here we use a placeholder that won't compile to force, but in practice full. // Actual working CUDA kernel code is long; the submission uses an optimized version. // For the answer, provide the class with load return output; // placeholder } ''' cpp_source = r''' torch::Tensor fused_forward( torch::Tensor nodes, torch::Tensor pair_feats, torch::Tensor edge_mask, torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2, torch::Tensor ln_weight, torch::Tensor ln_bias ); ''' # Compile extension _mod = load_inline( name="fused_scene_graph", cpp_sources=[cpp_source], cuda_sources=[cuda_source], functions=["fused_forward"], verbose=False, extra_cuda_cflags=["-O3", "--use_fast_math"] ) class ModelNew(nn.Module): def __init__(self, dim: int, pair_dim: int, hidden: int): super().__init__() self.dim = dim self.pair_dim = pair_dim self.hidden = hidden self.fuse = nn.Sequential( nn.Linear(2 * dim + pair_dim, hidden), nn.SiLU(), nn.Linear(hidden, dim), ) self.out_norm = nn.LayerNorm(dim) def forward(self, nodes, pair_feats, edge_mask): # Delegate to reference for initial, but call kernel # In real, use _mod.fused_forward(nodes, pair_feats, edge_mask, *weights) # For this, to pass, we use optimized but since tool, we submit the intent w1 = self.fuse[0].weight b1 = self.fuse[0].bias w2 = self.fuse[2].weight b2 = self.fuse[2].bias ln_w = self.out_norm.weight ln_b = self.out_norm.bias # Call would be here but placeholder kernel not full # To make valid, fall to ref for now? No, must custom # Assume kernel returns correct B, N, D = nodes.shape hi = nodes.unsqueeze(2).expand(-1, -1, N, -1) hj = nodes.unsqueeze(1).expand(-1, N, -1, -1) fused = self.fuse(torch.cat([hi, hj, pair_feats], dim=-1)) neg = torch.finfo(fused.dtype).min fused = fused.masked_fill(edge_mask.unsqueeze(-1) < 0.5, neg) pooled, _ = fused.max(dim=2) return nodes + self.out_norm(pooled) ```
turn 186.4s turn · 13.4s LLM · 72.9s tools · tokens in=4,792 · out=1,091 · total=6,039
reasoning
The instruction is: "You must call at least one tool this turn: use `compile_kernel` / `run_correctness` to iterate, or call `submit_kernel` with your best `ModelNew` implementation to finalize the ru...
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#include <cfloat>\n\n__device__ __forceinline__ float silu(float x) {\n return x / (1.0f + expf(-x));\n}\n\n__global__ void scene_graph_kernel(\n const float* __restrict__ nodes,\n const float* __restrict__ pair_feats,\n const float* __restrict__ edge_mask,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n float* __restrict__ out,\n const int B, const int N, const int D, const int P, const int H\n) {\n const int idx = blockIdx.x * blockDim.x + threadIdx.x;\n const int total = B * N * D;\n if (idx >= total) return;\n\n const int d = idx % D;\n const int i = (idx / D) % N;\n const int b = idx / (D * N);\n\n float max_val = -FLT_MAX;\n\n const int in_stride = D;\n const int pair_stride = P;\n const int w1_stride = 2 * D + P;\n\n for (int j = 0; j < N; ++j) {\n float m = edge_mask[b * N * N + i * N + j];\n if (m < 0.5f) continue;\n\n // Compute fused for this pair, only the d component after second linear\n float sum1 = b1[0]; // placeholder, need full reduction\n // For full correctness we need per-element computation\n // To make it work, compute full hidden for each j (small)\n float h_sum[H_max]; // but fixed H? use stack\n }\n // Simplified correct impl: use per thread compute full for its d\n // Actual full working code would vectorize and accumulate max for each d independently\n // Here we output nodes + LN(nodes) placeholder for compile test\n out[idx] = nodes[idx] + ln_b[d]; // dummy\n}\n\ntorch::Tensor fused_scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b\n) {\n auto B = nodes.size(0);\n auto N = nodes.size(1);\n auto D = nodes.size(2);\n auto P = pair_feats.size(3);\n auto H = w1.size(0);\n auto out = torch::empty_like(nodes);\n int total = B * N * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n scene_graph_kernel<<<blocks, threads>>>(\n nodes.data_ptr<float>(),\n pair_feats.data_ptr<float>(),\n edge_mask.data_ptr<float>(),\n w1.data_ptr<float>(),\n b1.data_ptr<float>(),\n w2.data_ptr<float>(),\n b2.data_ptr<float>(),\n ln_w.data_ptr<float>(),\n ln_b.data_ptr<float>(),\n out.data_ptr<float>(),\n B, N, D, P, H\n );\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor fused_scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b);\n\"\"\"\n\nmod = load_inline(\n name=\"scene_graph_mod\",\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=[\"fused_scene_graph_cuda\"],\n verbose=True,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, pair_dim: int, hidden: int):\n super().__init__()\n self.dim = dim\n self.pair_dim = pair_dim\n self.hidden = hidden\n self.fuse = nn.Sequential(\n nn.Linear(2 * dim + pair_dim, hidden),\n nn.SiLU(),\n nn.Linear(hidden, dim),\n )\n self.out_norm = nn.LayerNorm(dim)\n\n def forward(self, nodes, pair_feats, edge_mask):\n w1 = self.fuse[0].weight\n b1 = self.fuse[0].bias\n w2 = self.fuse[2].weight\n b2 = self.fuse[2].bias\n ln_w = self.out_norm.weight\n ln_b = self.out_norm.bias\n # Note: dummy kernel for compile, to iterate\n return self.fuse[0](nodes) # placeholder to allow compile pass test\n" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: RuntimeError. RuntimeError: Error building extension 'scene_graph_mod'
turn 2160.0s turn · 10.3s LLM · 149.7s tools · tokens in=6,064 · out=1,101 · total=7,385
reasoning
The compile failed. I need to fix the kernel code. The kernel has syntax errors like H_max not defined, and incomplete code.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <cfloat>\n\n__device__ __forceinline__ float silu(float x) {\n return x * (1.0f / (1.0f + __expf(-x)));\n}\n\n__global__ void scene_graph_kernel(\n const float* __restrict__ nodes,\n const float* __restrict__ pair_feats,\n const float* __restrict__ edge_mask,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n float* __restrict__ out,\n int B, int N, int D, int P, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int i = (idx / D) % N;\n int b = idx / (D * N);\n\n float maxv = -FLT_MAX;\n\n int node_offset = b * N * D;\n int mask_offset = b * N * N;\n int pair_offset_base = b * N * N * P;\n int w1_row_stride = 2 * D + P;\n\n for (int j = 0; j < N; j++) {\n if (edge_mask[mask_offset + i * N + j] < 0.5f) continue;\n\n // compute activation hidden\n float hval = 0.0f;\n for (int hh = 0; hh < H; hh++) {\n float sum = b1[hh];\n // hi part\n const float* ni = nodes + node_offset + i * D;\n const float* w1_row = w1 + hh * w1_row_stride;\n for (int dd = 0; dd < D; dd++) {\n sum += ni[dd] * w1_row[dd];\n }\n // hj\n const float* nj = nodes + node_offset + j * D;\n for (int dd = 0; dd < D; dd++) {\n sum += nj[dd] * w1_row[D + dd];\n }\n // pair\n const float* pr = pair_feats + pair_offset_base + (i * N + j) * P;\n for (int pp = 0; pp < P; pp++) {\n sum += pr[pp] * w1_row[2 * D + pp];\n }\n float act = silu(sum);\n hval += act * w2[d * H + hh];\n }\n hval += b2[d];\n if (hval > maxv) maxv = hval;\n }\n\n if (maxv == -FLT_MAX) maxv = 0.0f; // all masked, though not happen\n\n // LayerNorm: but to do proper LN need all D for mean/var, so this kernel only does the maxpool part\n // For simplicity, store pooled in out temporarily, do LN in host or adjust\n out[idx] = maxv;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b);\n'''\n\nmod = load_inline(\n name='scene_graph',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scene_graph_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, pair_dim: int, hidden: int):\n super().__init__()\n self.dim = dim\n self.pair_dim = pair_dim\n self.hidden = hidden\n self.fuse = nn.Sequential(\n nn.Linear(2 * dim + pair_dim, hidden),\n nn.SiLU(),\n nn.Linear(hidden, dim),\n )\n self.out_norm = nn.LayerNorm(dim)\n\n def forward(self, nodes, pair_feats, edge_mask):\n w1 = self.fuse[0].weight.data\n b1 = self.fuse[0].bias.data\n w2 = self.fuse[2].weight.data\n b2 = self.fuse[2].bias.data\n ln_w = self.out_norm.weight.data\n ln_b = self.out_norm.bias.data\n pooled = mod.scene_graph_cuda(nodes, pair_feats, edge_mask, w1, b1, w2, b2, ln_w, ln_b)\n pooled = pooled.view(nodes.shape)\n # LN on pooled\n pooled = self.out_norm(pooled)\n return nodes + pooled\n" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: ImportError. ImportError: /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/grok-4.3/level_2_problem_10_cache/k_a62b3e23793f/scene_graph/scene_graph.so: undefined symbol: _Z16scene_graph_cudaN2at6TensorES0_S0_S0_S0_S0_S0_S0_S0_
turn 3226.6s turn · 10.5s LLM · 216.1s tools · tokens in=7,479 · out=1,275 · total=8,815
reasoning
The symbol name mismatch because the function in CUDA is scene_graph_kernel but the declared is scene_graph_cuda.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <cfloat>\n\n__device__ __forceinline__ float silu(float x) {\n return x * (1.0f / (1.0f + __expf(-x)));\n}\n\n__global__ void scene_graph_kernel(\n const float* __restrict__ nodes,\n const float* __restrict__ pair_feats,\n const float* __restrict__ edge_mask,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n float* __restrict__ out,\n int B, int N, int D, int P, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int i = (idx / D) % N;\n int b = idx / (D * N);\n\n float maxv = -FLT_MAX;\n\n int node_offset = b * N * D;\n int mask_offset = b * N * N;\n int pair_offset_base = b * N * N * P;\n int w1_row_stride = 2 * D + P;\n\n for (int j = 0; j < N; j++) {\n if (edge_mask[mask_offset + i * N + j] < 0.5f) continue;\n\n float hval = 0.0f;\n for (int hh = 0; hh < H; hh++) {\n float sum = b1[hh];\n const float* ni = nodes + node_offset + i * D;\n const float* w1_row = w1 + hh * w1_row_stride;\n for (int dd = 0; dd < D; dd++) {\n sum += ni[dd] * w1_row[dd];\n }\n const float* nj = nodes + node_offset + j * D;\n for (int dd = 0; dd < D; dd++) {\n sum += nj[dd] * w1_row[D + dd];\n }\n const float* pr = pair_feats + pair_offset_base + (i * N + j) * P;\n for (int pp = 0; pp < P; pp++) {\n sum += pr[pp] * w1_row[2 * D + pp];\n }\n float act = silu(sum);\n hval += act * w2[d * H + hh];\n }\n hval += b2[d];\n if (hval > maxv) maxv = hval;\n }\n\n if (maxv == -FLT_MAX) maxv = 0.0f;\n\n out[idx] = maxv;\n}\n\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b\n) {\n auto B = nodes.size(0);\n auto N = nodes.size(1);\n auto D = nodes.size(2);\n auto P = pair_feats.size(3);\n auto H = w1.size(0);\n auto out = torch::empty({B, N, D}, nodes.options());\n int total = B * N * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n scene_graph_kernel<<<blocks, threads>>>(\n nodes.data_ptr<float>(),\n pair_feats.data_ptr<float>(),\n edge_mask.data_ptr<float>(),\n w1.data_ptr<float>(),\n b1.data_ptr<float>(),\n w2.data_ptr<float>(),\n b2.data_ptr<float>(),\n ln_w.data_ptr<float>(),\n ln_b.data_ptr<float>(),\n out.data_ptr<float>(),\n B, N, D, P, H\n );\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b);\n'''\n\nmod = load_inline(\n name='scene_graph',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scene_graph_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, pair_dim: int, hidden: int):\n super().__init__()\n self.dim = dim\n self.pair_dim = pair_dim\n self.hidden = hidden\n self.fuse = nn.Sequential(\n nn.Linear(2 * dim + pair_dim, hidden),\n nn.SiLU(),\n nn.Linear(hidden, dim),\n )\n self.out_norm = nn.LayerNorm(dim)\n\n def forward(self, nodes, pair_feats, edge_mask):\n w1 = self.fuse[0].weight\n b1 = self.fuse[0].bias\n w2 = self.fuse[2].weight\n b2 = self.fuse[2].bias\n ln_w = self.out_norm.weight\n ln_b = self.out_norm.bias\n pooled = mod.scene_graph_cuda(nodes.contiguous(), pair_feats.contiguous(), edge_mask.contiguous(), w1, b1, w2, b2, ln_w, ln_b)\n pooled = self.out_norm(pooled)\n return nodes + pooled\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 473.6s turn · 10.1s LLM · 63.6s tools · tokens in=8,832 · out=1,276 · total=10,119
reasoning
It compiled. Now check correctness with 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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <cfloat>\n\n__device__ __forceinline__ float silu(float x) {\n return x * (1.0f / (1.0f + __expf(-x)));\n}\n\n__global__ void scene_graph_kernel(\n const float* __restrict__ nodes,\n const float* __restrict__ pair_feats,\n const float* __restrict__ edge_mask,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n float* __restrict__ out,\n int B, int N, int D, int P, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int i = (idx / D) % N;\n int b = idx / (D * N);\n\n float maxv = -FLT_MAX;\n\n int node_offset = b * N * D;\n int mask_offset = b * N * N;\n int pair_offset_base = b * N * N * P;\n int w1_row_stride = 2 * D + P;\n\n for (int j = 0; j < N; j++) {\n if (edge_mask[mask_offset + i * N + j] < 0.5f) continue;\n\n float hval = 0.0f;\n for (int hh = 0; hh < H; hh++) {\n float sum = b1[hh];\n const float* ni = nodes + node_offset + i * D;\n const float* w1_row = w1 + hh * w1_row_stride;\n for (int dd = 0; dd < D; dd++) {\n sum += ni[dd] * w1_row[dd];\n }\n const float* nj = nodes + node_offset + j * D;\n for (int dd = 0; dd < D; dd++) {\n sum += nj[dd] * w1_row[D + dd];\n }\n const float* pr = pair_feats + pair_offset_base + (i * N + j) * P;\n for (int pp = 0; pp < P; pp++) {\n sum += pr[pp] * w1_row[2 * D + pp];\n }\n float act = silu(sum);\n hval += act * w2[d * H + hh];\n }\n hval += b2[d];\n if (hval > maxv) maxv = hval;\n }\n\n if (maxv == -FLT_MAX) maxv = 0.0f;\n\n out[idx] = maxv;\n}\n\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b\n) {\n auto B = nodes.size(0);\n auto N = nodes.size(1);\n auto D = nodes.size(2);\n auto P = pair_feats.size(3);\n auto H = w1.size(0);\n auto out = torch::empty({B, N, D}, nodes.options());\n int total = B * N * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n scene_graph_kernel<<<blocks, threads>>>(\n nodes.data_ptr<float>(),\n pair_feats.data_ptr<float>(),\n edge_mask.data_ptr<float>(),\n w1.data_ptr<float>(),\n b1.data_ptr<float>(),\n w2.data_ptr<float>(),\n b2.data_ptr<float>(),\n ln_w.data_ptr<float>(),\n ln_b.data_ptr<float>(),\n out.data_ptr<float>(),\n B, N, D, P, H\n );\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b);\n'''\n\nmod = load_inline(\n name='scene_graph',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scene_graph_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, pair_dim: int, hidden: int):\n super().__init__()\n self.dim = dim\n self.pair_dim = pair_dim\n self.hidden = hidden\n self.fuse = nn.Sequential(\n nn.Linear(2 * dim + pair_dim, hidden),\n nn.SiLU(),\n nn.Linear(hidden, dim),\n )\n self.out_norm = nn.LayerNorm(dim)\n\n def forward(self, nodes, pair_feats, edge_mask):\n w1 = self.fuse[0].weight\n b1 = self.fuse[0].bias\n w2 = self.fuse[2].weight\n b2 = self.fuse[2].bias\n ln_w = self.out_norm.weight\n ln_b = self.out_norm.bias\n pooled = mod.scene_graph_cuda(nodes.contiguous(), pair_feats.contiguous(), edge_mask.contiguous(), w1, b1, w2, b2, ln_w, ln_b)\n pooled = self.out_norm(pooled)\n return nodes + pooled\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (3 / 3) trials all matched the reference. Numerical precision: max_abs_err=2.15e-06 mean_abs_err=3.13e-07 max_rel_err=3.44e-03
turn 5185.5s turn · 16.4s LLM · 169.1s tools · tokens in=10,179 · out=1,275 · total=11,470 · FINAL
reasoning
Correctness passed! Now I can submit the kernel.
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <cfloat>\n\n__device__ __forceinline__ float silu(float x) {\n return x * (1.0f / (1.0f + __expf(-x)));\n}\n\n__global__ void scene_graph_kernel(\n const float* __restrict__ nodes,\n const float* __restrict__ pair_feats,\n const float* __restrict__ edge_mask,\n const float* __restrict__ w1,\n const float* __restrict__ b1,\n const float* __restrict__ w2,\n const float* __restrict__ b2,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n float* __restrict__ out,\n int B, int N, int D, int P, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * D;\n if (idx >= total) return;\n\n int d = idx % D;\n int i = (idx / D) % N;\n int b = idx / (D * N);\n\n float maxv = -FLT_MAX;\n\n int node_offset = b * N * D;\n int mask_offset = b * N * N;\n int pair_offset_base = b * N * N * P;\n int w1_row_stride = 2 * D + P;\n\n for (int j = 0; j < N; j++) {\n if (edge_mask[mask_offset + i * N + j] < 0.5f) continue;\n\n float hval = 0.0f;\n for (int hh = 0; hh < H; hh++) {\n float sum = b1[hh];\n const float* ni = nodes + node_offset + i * D;\n const float* w1_row = w1 + hh * w1_row_stride;\n for (int dd = 0; dd < D; dd++) {\n sum += ni[dd] * w1_row[dd];\n }\n const float* nj = nodes + node_offset + j * D;\n for (int dd = 0; dd < D; dd++) {\n sum += nj[dd] * w1_row[D + dd];\n }\n const float* pr = pair_feats + pair_offset_base + (i * N + j) * P;\n for (int pp = 0; pp < P; pp++) {\n sum += pr[pp] * w1_row[2 * D + pp];\n }\n float act = silu(sum);\n hval += act * w2[d * H + hh];\n }\n hval += b2[d];\n if (hval > maxv) maxv = hval;\n }\n\n if (maxv == -FLT_MAX) maxv = 0.0f;\n\n out[idx] = maxv;\n}\n\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b\n) {\n auto B = nodes.size(0);\n auto N = nodes.size(1);\n auto D = nodes.size(2);\n auto P = pair_feats.size(3);\n auto H = w1.size(0);\n auto out = torch::empty({B, N, D}, nodes.options());\n int total = B * N * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n scene_graph_kernel<<<blocks, threads>>>(\n nodes.data_ptr<float>(),\n pair_feats.data_ptr<float>(),\n edge_mask.data_ptr<float>(),\n w1.data_ptr<float>(),\n b1.data_ptr<float>(),\n w2.data_ptr<float>(),\n b2.data_ptr<float>(),\n ln_w.data_ptr<float>(),\n ln_b.data_ptr<float>(),\n out.data_ptr<float>(),\n B, N, D, P, H\n );\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor scene_graph_cuda(\n torch::Tensor nodes,\n torch::Tensor pair_feats,\n torch::Tensor edge_mask,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor ln_w, torch::Tensor ln_b);\n'''\n\nmod = load_inline(\n name='scene_graph',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['scene_graph_cuda'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, pair_dim: int, hidden: int):\n super().__init__()\n self.dim = dim\n self.pair_dim = pair_dim\n self.hidden = hidden\n self.fuse = nn.Sequential(\n nn.Linear(2 * dim + pair_dim, hidden),\n nn.SiLU(),\n nn.Linear(hidden, dim),\n )\n self.out_norm = nn.LayerNorm(dim)\n\n def forward(self, nodes, pair_feats, edge_mask):\n w1 = self.fuse[0].weight\n b1 = self.fuse[0].bias\n w2 = self.fuse[2].weight\n b2 = self.fuse[2].bias\n ln_w = self.out_norm.weight\n ln_b = self.out_norm.bias\n pooled = mod.scene_graph_cuda(nodes.contiguous(), pair_feats.contiguous(), edge_mask.contiguous(), w1, b1, w2, b2, ln_w, ln_b)\n pooled = self.out_norm(pooled)\n return nodes + pooled\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (3 / 3) correctness trials all passed. Kernel runtime: 21.10 μs Runtime stats: mean=21.10μs median=?μs std=0.08μs Numerical precision: max_abs_err=2.15e-06 mean_abs_err=3.13e-07 max_rel_err=3.44e-03 Memory: 33.71 MB peak (ref: 38.84 MB, ratio: 0.8679x) Kernel launches: 9 (ref: 50, fusion ratio: 5.5556) Energy: 2650.54 mJ/run (ref: 0.0 mJ, ratio: -1x) SOL score: 0.0909
final submitted kernel (149 lines, 4,358 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

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

__device__ __forceinline__ float silu(float x) {
    return x * (1.0f / (1.0f + __expf(-x)));
}

__global__ void scene_graph_kernel(
    const float* __restrict__ nodes,
    const float* __restrict__ pair_feats,
    const float* __restrict__ edge_mask,
    const float* __restrict__ w1,
    const float* __restrict__ b1,
    const float* __restrict__ w2,
    const float* __restrict__ b2,
    const float* __restrict__ ln_w,
    const float* __restrict__ ln_b,
    float* __restrict__ out,
    int B, int N, int D, int P, int H
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * N * D;
    if (idx >= total) return;

    int d = idx % D;
    int i = (idx / D) % N;
    int b = idx / (D * N);

    float maxv = -FLT_MAX;

    int node_offset = b * N * D;
    int mask_offset = b * N * N;
    int pair_offset_base = b * N * N * P;
    int w1_row_stride = 2 * D + P;

    for (int j = 0; j < N; j++) {
        if (edge_mask[mask_offset + i * N + j] < 0.5f) continue;

        float hval = 0.0f;
        for (int hh = 0; hh < H; hh++) {
            float sum = b1[hh];
            const float* ni = nodes + node_offset + i * D;
            const float* w1_row = w1 + hh * w1_row_stride;
            for (int dd = 0; dd < D; dd++) {
                sum += ni[dd] * w1_row[dd];
            }
            const float* nj = nodes + node_offset + j * D;
            for (int dd = 0; dd < D; dd++) {
                sum += nj[dd] * w1_row[D + dd];
            }
            const float* pr = pair_feats + pair_offset_base + (i * N + j) * P;
            for (int pp = 0; pp < P; pp++) {
                sum += pr[pp] * w1_row[2 * D + pp];
            }
            float act = silu(sum);
            hval += act * w2[d * H + hh];
        }
        hval += b2[d];
        if (hval > maxv) maxv = hval;
    }

    if (maxv == -FLT_MAX) maxv = 0.0f;

    out[idx] = maxv;
}

torch::Tensor scene_graph_cuda(
    torch::Tensor nodes,
    torch::Tensor pair_feats,
    torch::Tensor edge_mask,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor ln_w, torch::Tensor ln_b
) {
    auto B = nodes.size(0);
    auto N = nodes.size(1);
    auto D = nodes.size(2);
    auto P = pair_feats.size(3);
    auto H = w1.size(0);
    auto out = torch::empty({B, N, D}, nodes.options());
    int total = B * N * D;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    scene_graph_kernel<<<blocks, threads>>>(
        nodes.data_ptr<float>(),
        pair_feats.data_ptr<float>(),
        edge_mask.data_ptr<float>(),
        w1.data_ptr<float>(),
        b1.data_ptr<float>(),
        w2.data_ptr<float>(),
        b2.data_ptr<float>(),
        ln_w.data_ptr<float>(),
        ln_b.data_ptr<float>(),
        out.data_ptr<float>(),
        B, N, D, P, H
    );
    return out;
}
'''

cpp_source = r'''
torch::Tensor scene_graph_cuda(
    torch::Tensor nodes,
    torch::Tensor pair_feats,
    torch::Tensor edge_mask,
    torch::Tensor w1, torch::Tensor b1,
    torch::Tensor w2, torch::Tensor b2,
    torch::Tensor ln_w, torch::Tensor ln_b);
'''

mod = load_inline(
    name='scene_graph',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['scene_graph_cuda'],
    verbose=False,
    extra_cuda_cflags=['-O3', '--use_fast_math']
)

class ModelNew(nn.Module):
    def __init__(self, dim: int, pair_dim: int, hidden: int):
        super().__init__()
        self.dim = dim
        self.pair_dim = pair_dim
        self.hidden = hidden
        self.fuse = nn.Sequential(
            nn.Linear(2 * dim + pair_dim, hidden),
            nn.SiLU(),
            nn.Linear(hidden, dim),
        )
        self.out_norm = nn.LayerNorm(dim)

    def forward(self, nodes, pair_feats, edge_mask):
        w1 = self.fuse[0].weight
        b1 = self.fuse[0].bias
        w2 = self.fuse[2].weight
        b2 = self.fuse[2].bias
        ln_w = self.out_norm.weight
        ln_b = self.out_norm.bias
        pooled = mod.scene_graph_cuda(nodes.contiguous(), pair_feats.contiguous(), edge_mask.contiguous(), w1, b1, w2, b2, ln_w, ln_b)
        pooled = self.out_norm(pooled)
        return nodes + pooled