correct
10_SampledDenseDenseMatmulEdges.py
started
2026-05-11T20:03:11Z
timingkernelbench.timing
kernel mean / median / std
0.05 / — / 3.24e-04 μs
ref mean / median / std
0.05 / — / 4.31e-04 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
2.11x
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
class Model(nn.Module):
"""
Edge-only sampled dense-dense matmul.
For each graph edge (src, dst), compute a dot product between source and
destination node embeddings. This pattern appears in link prediction,
sampled attention, and graph contrastive learning.
"""
def __init__(self):
super().__init__()
def forward(
self,
src_idx: torch.Tensor,
dst_idx: torch.Tensor,
src_feat: torch.Tensor,
dst_feat: torch.Tensor,
) -> torch.Tensor:
src = src_feat[src_idx.long()]
dst = dst_feat[dst_idx.long()]
return (src * dst).sum(dim=-1)
num_src_nodes = 1024
num_dst_nodes = 768
num_edges = 16384
feat_dim = 128
def get_inputs():
p = popcorn_pri
n_edges = p.jitter_int(num_edges)
fd = p.jitter_int(feat_dim, align=8)
ns = p.jitter_int(num_src_nodes)
nd = p.jitter_int(num_dst_nodes)
src_idx = torch.randint(0, ns, (n_edges,), dtype=torch.int32)
dst_idx = torch.randint(0, nd, (n_edges,), dtype=torch.int32)
src_feat = torch.randn(ns, fd, dtype=torch.float32)
dst_feat = torch.randn(nd, fd, dtype=torch.float32)
return [src_idx, dst_idx, src_feat, dst_feat]
def get_init_inputs():
return []
turn 0128.0s turn · 12.5s LLM · 115.6s tools · tokens in=2,052 · out=498 · total=3,407
reasoning
The task is to replace the PyTorch model's forward with a custom CUDA kernel for an edge-only sampled dense-dense matmul. It computes dot products between source and destination node embeddings for ea...The task involves replacing the forward pass of a PyTorch model with a custom CUDA kernel for computing dot products between node embeddings for graph edges. The reference PyTorch code gathers source and destination features using indices and computes their dot product per edge. For ModelNew, I need to accept the same inputs, output the same tensors tensors: src_idx, dst_idx, src_feat, dst_feat, and internally use a CUDA kernel to speed up the computation.
To implement this, I'll use torch.utils.cpp_extension.load_inline to compile the CUDA code inline. The kernel will launch threads to handle edges in parallel.
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\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n\nmodule = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 193.3s turn · 4.4s LLM · 88.8s tools · tokens in=3,424 · out=499 · total=3,930
reasoning
Compiled successfully. Now check correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n\nmodule = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Int but found Float
turn 291.7s turn · 15.0s LLM · 76.7s tools · tokens in=3,959 · out=534 · total=4,845
reasoning
Error: expected scalar type Int but found Float. The indices are int32, but in PyTorch, indexing expects long? No, the error is probably in data_ptr or something.
I encountered an error saying "expected scalar type Int but found Float" when replacing the forward computation with a CUDA kernel. The indices are torch.int32, but PyTorch indexing expects long tensors, which explains the casting needed in the reference implementation.
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\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n\nmodule = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 37.5s turn · 7.5s LLM · 0.0s tools · tokens in=4,862 · out=535 · total=5,402
reasoning
Compiled. Now 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\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n\nmodule = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Int but found Float
turn 413.8s turn · 13.8s LLM · 0.0s tools · tokens in=5,431 · out=533 · total=6,498
reasoning
Still the same error. Hmm.
The thinking trace shows that despite repeated attempts, the same error "expected scalar type Int but found Float" persists. This PyTorch error indicates a type mismatch, possibly when using a float tensor instead of an int one, like in indexing. The reference implementation uses src_idx.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 585.2s turn · 7.0s LLM · 78.2s tools · tokens in=6,515 · out=534 · total=7,053
reasoning
Now test correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n\n__global__ void edge_dot_product(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Int but found Float
turn 68.1s turn · 8.1s LLM · 0.0s tools · tokens in=7,082 · out=550 · total=7,820
reasoning
Still failing. The error must be coming from the reference implementation when run_correctness compares.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cstdint>\n\n__global__ void edge_dot_product(\n const int32_t* __restrict__ src_idx,\n const int32_t* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int32_t src = src_idx[edge];\n int32_t dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int32_t>(),\n dst_idx.data_ptr<int32_t>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 781.7s turn · 5.9s LLM · 75.8s tools · tokens in=7,837 · out=551 · total=8,392
reasoning
Try correctness again.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cstdint>\n\n__global__ void edge_dot_product(\n const int32_t* __restrict__ src_idx,\n const int32_t* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int32_t src = src_idx[edge];\n int32_t dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous();\n dst_idx = dst_idx.contiguous();\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int32_t>(),\n dst_idx.data_ptr<int32_t>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Int but found Float
turn 87.3s turn · 7.3s LLM · 0.0s tools · tokens in=8,421 · out=564 · total=9,197
reasoning
Still same. Perhaps the test is failing because tensors are not CUDA.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cstdint>\n\n__global__ void edge_dot_product(\n const int32_t* __restrict__ src_idx,\n const int32_t* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int32_t src = src_idx[edge];\n int32_t dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous().to(torch::kInt32);\n dst_idx = dst_idx.contiguous().to(torch::kInt32);\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int32_t>(),\n dst_idx.data_ptr<int32_t>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 997.2s turn · 6.4s LLM · 90.8s tools · tokens in=9,214 · out=565 · total=9,781
reasoning
Test.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cstdint>\n\n__global__ void edge_dot_product(\n const int32_t* __restrict__ src_idx,\n const int32_t* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int32_t src = src_idx[edge];\n int32_t dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous().to(torch::kInt32);\n dst_idx = dst_idx.contiguous().to(torch::kInt32);\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int32_t>(),\n dst_idx.data_ptr<int32_t>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=1.91e-05 mean_abs_err=1.75e-06 max_rel_err=4.33e-03
turn 10162.3s turn · 6.6s LLM · 155.6s tools · tokens in=9,841 · out=564 · total=10,453 · FINAL
reasoning
Great! Now it's correct. So the to(torch::kInt32) fixed the type issue somehow, probably inputs can be long or whatever.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n cuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cstdint>\n\n__global__ void edge_dot_product(\n const int32_t* __restrict__ src_idx,\n const int32_t* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n int edge = blockIdx.x * blockDim.x + threadIdx.x;\n if (edge < num_edges) {\n int32_t src = src_idx[edge];\n int32_t dst = dst_idx[edge];\n float sum = 0.0f;\n for (int d = 0; d < feat_dim; d++) {\n sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];\n }\n output[edge] = sum;\n }\n}\n\ntorch::Tensor edge_dot(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n src_idx = src_idx.contiguous().to(torch::kInt32);\n dst_idx = dst_idx.contiguous().to(torch::kInt32);\n src_feat = src_feat.contiguous();\n dst_feat = dst_feat.contiguous();\n int num_edges = src_idx.size(0);\n int feat_dim = src_feat.size(1);\n auto output = torch::empty({num_edges}, src_feat.options());\n dim3 block(256);\n dim3 grid((num_edges + 255) / 256);\n edge_dot_product<<<grid, block>>>(\n src_idx.data_ptr<int32_t>(),\n dst_idx.data_ptr<int32_t>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n return output;\n}\n\"\"\"\n cpp_source = \"\"\"\ntorch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);\n\"\"\"\n self.module = load_inline(\n name='edge_dot_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['edge_dot'],\n verbose=False\n )\n\n def forward(self, src_idx, dst_idx, src_feat, dst_feat):\n return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.05 μs
Runtime stats: mean=0.05μs median=?μs std=0.00μs
Numerical precision: max_abs_err=1.91e-05 mean_abs_err=1.75e-06 max_rel_err=4.33e-03
Memory: 34.19 MB peak (ref: 58.06 MB, ratio: 0.5888x)
Kernel launches: 9 (ref: 19, fusion ratio: 2.1111)
SOL score: 0.02
final submitted kernel (73 lines, 2,173 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
cuda_source = """
#include <torch/extension.h>
#include <cstdint>
__global__ void edge_dot_product(
const int32_t* __restrict__ src_idx,
const int32_t* __restrict__ dst_idx,
const float* __restrict__ src_feat,
const float* __restrict__ dst_feat,
float* __restrict__ output,
int num_edges,
int feat_dim
) {
int edge = blockIdx.x * blockDim.x + threadIdx.x;
if (edge < num_edges) {
int32_t src = src_idx[edge];
int32_t dst = dst_idx[edge];
float sum = 0.0f;
for (int d = 0; d < feat_dim; d++) {
sum += src_feat[src * feat_dim + d] * dst_feat[dst * feat_dim + d];
}
output[edge] = sum;
}
}
torch::Tensor edge_dot(
torch::Tensor src_idx,
torch::Tensor dst_idx,
torch::Tensor src_feat,
torch::Tensor dst_feat
) {
src_idx = src_idx.contiguous().to(torch::kInt32);
dst_idx = dst_idx.contiguous().to(torch::kInt32);
src_feat = src_feat.contiguous();
dst_feat = dst_feat.contiguous();
int num_edges = src_idx.size(0);
int feat_dim = src_feat.size(1);
auto output = torch::empty({num_edges}, src_feat.options());
dim3 block(256);
dim3 grid((num_edges + 255) / 256);
edge_dot_product<<<grid, block>>>(
src_idx.data_ptr<int32_t>(),
dst_idx.data_ptr<int32_t>(),
src_feat.data_ptr<float>(),
dst_feat.data_ptr<float>(),
output.data_ptr<float>(),
num_edges,
feat_dim
);
return output;
}
"""
cpp_source = """
torch::Tensor edge_dot(torch::Tensor src_idx, torch::Tensor dst_idx, torch::Tensor src_feat, torch::Tensor dst_feat);
"""
self.module = load_inline(
name='edge_dot_module',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['edge_dot'],
verbose=False
)
def forward(self, src_idx, dst_idx, src_feat, dst_feat):
return self.module.edge_dot(src_idx, dst_idx, src_feat, dst_feat)