correct
21_TriangularMultiplicativeUpdateIncoming.py
started
2026-05-11T21:08:14Z
timingkernelbench.timing
kernel mean / median / std
0.29 / — / 0.09 μs
ref mean / median / std
0.34 / — / 0.13 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
1.44x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
1749.32 GB/s
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
class Model(nn.Module):
"""
Triangular multiplicative update – incoming edges.
Mirror of the outgoing variant: updates pair (i,j) by aggregating over
the shared index k via edges (k,i) and (k,j).
"""
def __init__(self, pair_dim, hidden_dim):
super().__init__()
self.layer_norm_in = nn.LayerNorm(pair_dim)
self.proj_a = nn.Linear(pair_dim, hidden_dim)
self.gate_a = nn.Linear(pair_dim, hidden_dim)
self.proj_b = nn.Linear(pair_dim, hidden_dim)
self.gate_b = nn.Linear(pair_dim, hidden_dim)
self.layer_norm_out = nn.LayerNorm(hidden_dim)
self.out_proj = nn.Linear(hidden_dim, pair_dim)
self.out_gate = nn.Linear(pair_dim, pair_dim)
def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:
"""
Args:
pair_repr: (B, N, N, pair_dim)
pair_mask: (B, N, N)
Returns:
updated pair representation (B, N, N, pair_dim)
"""
mask = pair_mask.unsqueeze(-1)
x = self.layer_norm_in(pair_repr)
a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask # (B, N, N, H)
b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask
# Incoming: aggregate over k with a[k,i] * b[k,j]
out = torch.einsum("bkih,bkjh->bijh", a, b)
out = self.layer_norm_out(out)
out = self.out_proj(out)
out = torch.sigmoid(self.out_gate(pair_repr)) * out
return out
pair_dim = 64
hidden_dim = 32
seq_len = 32
batch_size = 2
def get_inputs():
p = popcorn_pri
pair_repr = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(seq_len), pair_dim)
pair_mask = torch.ones(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(seq_len))
return [pair_repr, pair_mask]
def get_init_inputs():
return [pair_dim, hidden_dim]
turn 0220.5s turn · 8.8s LLM · 211.6s tools · tokens in=2,254 · out=6 · total=2,937
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a "Triangular multiplicative update – incoming edges."
I need to create a ModelNew class that replaces the forward computation of the original PyTorch model with a custom CUDA kernel, matching the constructor args and forward signature, and producing the same numerical outputs. The task cuts off at "Produce numerically ...(2450 chars truncated)...ize = 2", but I can infer it's batch_size=2, probably for testing.
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 1221.2s turn · 12.3s LLM · 208.9s tools · tokens in=3,170 · out=749 · total=4,210
reasoning
H100, high bandwidth.
Now, I need to implement.
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 incoming_triangular_mul(\n const float* __restrict__ a,\n const float* __restrict__ b,\n float* __restrict__ out,\n int B, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H;\n if (idx >= total) return;\n\n int h = idx % H;\n int j = (idx / H) % N;\n int i = (idx / (H * N)) % N;\n int batch = idx / (H * N * N);\n\n float sum = 0.0f;\n for (int k = 0; k < N; ++k) {\n float aval = a[((batch * N + k) * N + i) * H + h];\n float bval = b[((batch * N + k) * N + j) * H + h];\n sum += aval * bval;\n }\n out[((batch * N + i) * N + j) * H + h] = sum;\n}\n\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b) {\n auto B = a.size(0);\n auto N = a.size(1);\n auto H = a.size(3);\n auto out = torch::zeros({B, N, N, H}, a.options());\n int total = B * N * N * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n incoming_triangular_mul<<<blocks, threads>>>(\n a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(),\n B, N, H\n );\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b);\n\"\"\"\n\nmodule = load_inline(\n name='incoming_mul',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['incoming_mul_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm_in = nn.LayerNorm(pair_dim)\n self.proj_a = nn.Linear(pair_dim, hidden_dim)\n self.gate_a = nn.Linear(pair_dim, hidden_dim)\n self.proj_b = nn.Linear(pair_dim, hidden_dim)\n self.gate_b = nn.Linear(pair_dim, hidden_dim)\n self.layer_norm_out = nn.LayerNorm(hidden_dim)\n self.out_proj = nn.Linear(hidden_dim, pair_dim)\n self.out_gate = nn.Linear(pair_dim, pair_dim)\n self.mul = module\n\n def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:\n mask = pair_mask.unsqueeze(-1)\n x = self.layer_norm_in(pair_repr)\n a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask\n b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask\n out = self.mul.incoming_mul_cuda(a, b)\n out = self.layer_norm_out(out)\n out = self.out_proj(out)\n out = torch.sigmoid(self.out_gate(pair_repr)) * out\n return out\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2149.8s turn · 7.7s LLM · 142.1s tools · tokens in=4,227 · out=750 · total=4,984
reasoning
Compiled ok. 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 incoming_triangular_mul(\n const float* __restrict__ a,\n const float* __restrict__ b,\n float* __restrict__ out,\n int B, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H;\n if (idx >= total) return;\n\n int h = idx % H;\n int j = (idx / H) % N;\n int i = (idx / (H * N)) % N;\n int batch = idx / (H * N * N);\n\n float sum = 0.0f;\n for (int k = 0; k < N; ++k) {\n float aval = a[((batch * N + k) * N + i) * H + h];\n float bval = b[((batch * N + k) * N + j) * H + h];\n sum += aval * bval;\n }\n out[((batch * N + i) * N + j) * H + h] = sum;\n}\n\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b) {\n auto B = a.size(0);\n auto N = a.size(1);\n auto H = a.size(3);\n auto out = torch::zeros({B, N, N, H}, a.options());\n int total = B * N * N * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n incoming_triangular_mul<<<blocks, threads>>>(\n a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(),\n B, N, H\n );\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b);\n\"\"\"\n\nmodule = load_inline(\n name='incoming_mul',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['incoming_mul_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm_in = nn.LayerNorm(pair_dim)\n self.proj_a = nn.Linear(pair_dim, hidden_dim)\n self.gate_a = nn.Linear(pair_dim, hidden_dim)\n self.proj_b = nn.Linear(pair_dim, hidden_dim)\n self.gate_b = nn.Linear(pair_dim, hidden_dim)\n self.layer_norm_out = nn.LayerNorm(hidden_dim)\n self.out_proj = nn.Linear(hidden_dim, pair_dim)\n self.out_gate = nn.Linear(pair_dim, pair_dim)\n self.mul = module\n\n def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:\n mask = pair_mask.unsqueeze(-1)\n x = self.layer_norm_in(pair_repr)\n a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask\n b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask\n out = self.mul.incoming_mul_cuda(a, b)\n out = self.layer_norm_out(out)\n out = self.out_proj(out)\n out = torch.sigmoid(self.out_gate(pair_repr)) * out\n return out\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 3126.3s turn · 8.4s LLM · 117.9s tools · tokens in=5,044 · out=749 · total=5,817
reasoning
Correct! Now, to check if faster, I need to profile or submit, but since submit ends, let's profile.
call → profile_kernel
callprofile_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 incoming_triangular_mul(\n const float* __restrict__ a,\n const float* __restrict__ b,\n float* __restrict__ out,\n int B, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H;\n if (idx >= total) return;\n\n int h = idx % H;\n int j = (idx / H) % N;\n int i = (idx / (H * N)) % N;\n int batch = idx / (H * N * N);\n\n float sum = 0.0f;\n for (int k = 0; k < N; ++k) {\n float aval = a[((batch * N + k) * N + i) * H + h];\n float bval = b[((batch * N + k) * N + j) * H + h];\n sum += aval * bval;\n }\n out[((batch * N + i) * N + j) * H + h] = sum;\n}\n\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b) {\n auto B = a.size(0);\n auto N = a.size(1);\n auto H = a.size(3);\n auto out = torch::zeros({B, N, N, H}, a.options());\n int total = B * N * N * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n incoming_triangular_mul<<<blocks, threads>>>(\n a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(),\n B, N, H\n );\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b);\n\"\"\"\n\nmodule = load_inline(\n name='incoming_mul',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['incoming_mul_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm_in = nn.LayerNorm(pair_dim)\n self.proj_a = nn.Linear(pair_dim, hidden_dim)\n self.gate_a = nn.Linear(pair_dim, hidden_dim)\n self.proj_b = nn.Linear(pair_dim, hidden_dim)\n self.gate_b = nn.Linear(pair_dim, hidden_dim)\n self.layer_norm_out = nn.LayerNorm(hidden_dim)\n self.out_proj = nn.Linear(hidden_dim, pair_dim)\n self.out_gate = nn.Linear(pair_dim, pair_dim)\n self.mul = module\n\n def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:\n mask = pair_mask.unsqueeze(-1)\n x = self.layer_norm_in(pair_repr)\n a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask\n b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask\n out = self.mul.incoming_mul_cuda(a, b)\n out = self.layer_norm_out(out)\n out = self.out_proj(out)\n out = torch.sigmoid(self.out_gate(pair_repr)) * out\n return out\n"
}
result ← profile_kernel (ok)
okprofile_kernel
profile_kernel PASSED: profiling complete.
=== Kernel Profile Summary ===
--- Kernel Breakdown (by GPU time) ---
#1: aten::linear 71551.0μs (39%) calls=1
#2: aten::addmm 71397.0μs (39%) calls=1
#3: aten::layer_norm 7584.0μs (4%) calls=1
#4: aten::native_layer_norm 7555.0μs (4%) calls=1
#5: aten::mul 5575.0μs (3%) calls=1
... +65 more kernels
--- Timing ---
GPU time: 97.79 μs
--- Memory Bandwidth ---
DRAM BW: 1688.5/3350 GB/s (50.4%)
read=1688.5 write=0.0 GB/s
Hit rates: L2=1468.0%
--- Compute Throughput ---
Dominant pipe: tensor-core util=6.9%
Pipe util: FMA=79.2%, Tensor=6.9%
--- Roofline Analysis ---
Ridge point (tensor-core): 295.22 FLOPs/byte
Classification: MEMORY-BOUND
--- Occupancy ---
Achieved: 213.5%
Eligible warps/cycle: 4.30
Launch: regs/thread=834, smem/blk=60160B, block=2048, grid=5952
--- Memory Access Quality ---
Global ld coalescing: 141.00 sectors/req (141.0x wasted BW)
--- Warp Stall Reasons (% of issue cycles) ---
long_scoreboard 18313.6% ####################
short_scoreboard 2468.3% ####################
mio_throttle 273.0% ####################
--- Optimization Hints ---
-> 18314% stalls: global memory latency (waiting for DRAM/L2). + sectors/req=141.0 -> uncoalesced global loads, restructure access pattern or use __ldg Try shared-memory tiling, prefetching, or reducing working set.
-> Global load coalescing=141.0x (ideal=1.0) -> 141x wasted DRAM BW. Make per-warp accesses contiguous.
turn 4259.0s turn · 9.1s LLM · 249.8s tools · tokens in=6,321 · out=749 · total=7,160 · FINAL
reasoning
It's faster? The GPU time is 97us for my kernel, but reference would have more.
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 incoming_triangular_mul(\n const float* __restrict__ a,\n const float* __restrict__ b,\n float* __restrict__ out,\n int B, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H;\n if (idx >= total) return;\n\n int h = idx % H;\n int j = (idx / H) % N;\n int i = (idx / (H * N)) % N;\n int batch = idx / (H * N * N);\n\n float sum = 0.0f;\n for (int k = 0; k < N; ++k) {\n float aval = a[((batch * N + k) * N + i) * H + h];\n float bval = b[((batch * N + k) * N + j) * H + h];\n sum += aval * bval;\n }\n out[((batch * N + i) * N + j) * H + h] = sum;\n}\n\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b) {\n auto B = a.size(0);\n auto N = a.size(1);\n auto H = a.size(3);\n auto out = torch::zeros({B, N, N, H}, a.options());\n int total = B * N * N * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n incoming_triangular_mul<<<blocks, threads>>>(\n a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(),\n B, N, H\n );\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b);\n\"\"\"\n\nmodule = load_inline(\n name='incoming_mul',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['incoming_mul_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm_in = nn.LayerNorm(pair_dim)\n self.proj_a = nn.Linear(pair_dim, hidden_dim)\n self.gate_a = nn.Linear(pair_dim, hidden_dim)\n self.proj_b = nn.Linear(pair_dim, hidden_dim)\n self.gate_b = nn.Linear(pair_dim, hidden_dim)\n self.layer_norm_out = nn.LayerNorm(hidden_dim)\n self.out_proj = nn.Linear(hidden_dim, pair_dim)\n self.out_gate = nn.Linear(pair_dim, pair_dim)\n self.mul = module\n\n def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:\n mask = pair_mask.unsqueeze(-1)\n x = self.layer_norm_in(pair_repr)\n a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask\n b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask\n out = self.mul.incoming_mul_cuda(a, b)\n out = self.layer_norm_out(out)\n out = self.out_proj(out)\n out = torch.sigmoid(self.out_gate(pair_repr)) * out\n return out\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.29 μs
Runtime stats: mean=0.29μs median=?μs std=0.09μs
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
Memory: 36.13 MB peak (ref: 36.13 MB, ratio: 1.0x)
Kernel launches: 70 (ref: 101, fusion ratio: 1.4429)
SOL score: 0.5222
final submitted kernel (85 lines, 2,619 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 incoming_triangular_mul(
const float* __restrict__ a,
const float* __restrict__ b,
float* __restrict__ out,
int B, int N, int H
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * N * N * H;
if (idx >= total) return;
int h = idx % H;
int j = (idx / H) % N;
int i = (idx / (H * N)) % N;
int batch = idx / (H * N * N);
float sum = 0.0f;
for (int k = 0; k < N; ++k) {
float aval = a[((batch * N + k) * N + i) * H + h];
float bval = b[((batch * N + k) * N + j) * H + h];
sum += aval * bval;
}
out[((batch * N + i) * N + j) * H + h] = sum;
}
torch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b) {
auto B = a.size(0);
auto N = a.size(1);
auto H = a.size(3);
auto out = torch::zeros({B, N, N, H}, a.options());
int total = B * N * N * H;
int threads = 256;
int blocks = (total + threads - 1) / threads;
incoming_triangular_mul<<<blocks, threads>>>(
a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(),
B, N, H
);
return out;
}
"""
cpp_source = """
torch::Tensor incoming_mul_cuda(torch::Tensor a, torch::Tensor b);
"""
module = load_inline(
name='incoming_mul',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['incoming_mul_cuda'],
verbose=True
)
class ModelNew(nn.Module):
def __init__(self, pair_dim, hidden_dim):
super().__init__()
self.layer_norm_in = nn.LayerNorm(pair_dim)
self.proj_a = nn.Linear(pair_dim, hidden_dim)
self.gate_a = nn.Linear(pair_dim, hidden_dim)
self.proj_b = nn.Linear(pair_dim, hidden_dim)
self.gate_b = nn.Linear(pair_dim, hidden_dim)
self.layer_norm_out = nn.LayerNorm(hidden_dim)
self.out_proj = nn.Linear(hidden_dim, pair_dim)
self.out_gate = nn.Linear(pair_dim, pair_dim)
self.mul = module
def forward(self, pair_repr: torch.Tensor, pair_mask: torch.Tensor) -> torch.Tensor:
mask = pair_mask.unsqueeze(-1)
x = self.layer_norm_in(pair_repr)
a = self.proj_a(x) * torch.sigmoid(self.gate_a(x)) * mask
b = self.proj_b(x) * torch.sigmoid(self.gate_b(x)) * mask
out = self.mul.incoming_mul_cuda(a, b)
out = self.layer_norm_out(out)
out = self.out_proj(out)
out = torch.sigmoid(self.out_gate(pair_repr)) * out
return out