correct
10_GFVLA_FusedPairwiseNodeAggregate.py
started
2026-05-11T17:36:20Z
timingkernelbench.timing
kernel mean / median / std
21.10 / — / 0.08 μs
ref mean / median / std
0.16 / — / 0.05 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
5.56x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
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