incorrect
17_CSRMaxAggregation.py
started
2026-05-11T12:06:57Z
correctness issue:
timingkernelbench.timing
kernel mean / median / std
—
ref mean / median / std
1550.00 / — / 0.37 μs
roofline / occupancynsight_hardware_counters
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Per-destination max aggregation over neighbor features in CSR format.
This pattern appears in max-pooling GNNs and neighborhood feature
extraction where reductions are not sums but feature-wise maxima.
"""
def __init__(self):
super().__init__()
def forward(
self,
row_ptr: torch.Tensor,
col_idx: torch.Tensor,
node_feat: torch.Tensor,
) -> torch.Tensor:
num_nodes = row_ptr.numel() - 1
feat_dim = node_feat.shape[1]
out = torch.full((num_nodes, feat_dim), float("-inf"), dtype=node_feat.dtype, device=node_feat.device)
for dst in range(num_nodes):
start = int(row_ptr[dst].item())
end = int(row_ptr[dst + 1].item())
if end > start:
out[dst] = node_feat[col_idx[start:end].long()].amax(dim=0)
return out
num_nodes = 640
avg_degree = 16
feat_dim = 80
def get_inputs():
return list(popcorn_pri.make_csr_graph(num_nodes, avg_degree, feat_dim))
def get_init_inputs():
return []
turn 0377.1s turn · 15.8s LLM · 361.3s tools · tokens in=1,626 · out=6 · total=2,051
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs per-destination max aggregation over neighbor features in CSR format.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs a per-destination max aggregation on neighbor features in CSR format. The inputs are row_ptr for CSR row pointers, col_idx for column indices, and node_feat for node features.
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 1614.6s turn · 11.9s LLM · 602.6s tools · tokens in=2,284 · out=611 · total=2,988
reasoning
H100, high bandwidth, but for this small problem, it's about kernel launch efficiency.
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#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = -INFINITY;\n }\n return;\n }\n\n extern __shared__ float sdata[];\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = -INFINITY;\n for (int i = 0; i < degree; ++i) {\n int src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n size_t shared_mem = 0; // not using shared for now\n\n max_aggregate_kernel<<<blocks, threads, shared_mem>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2336.1s turn · 7.7s LLM · 328.4s tools · tokens in=3,005 · out=596 · total=3,608
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#include <cuda.h>\n#include <cuda_runtime.h>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = -INFINITY;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = -INFINITY;\n for (int i = 0; i < degree; ++i) {\n int src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n size_t shared_mem = 0;\n\n max_aggregate_kernel<<<blocks, threads, shared_mem>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_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 3503.1s turn · 8.2s LLM · 494.9s tools · tokens in=3,637 · out=598 · total=4,421
reasoning
Error: expected scalar type Int but found Float. Probably row_ptr or col_idx are int32, but in C++ it's int which may be 64? No.
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#include <cstdint>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = -INFINITY;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = -INFINITY;\n for (int i = 0; i < degree; ++i) {\n int32_t src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n\n max_aggregate_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 4430.0s turn · 6.1s LLM · 423.9s tools · tokens in=4,438 · out=599 · total=5,041
reasoning
Now 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 = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cstdint>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = -INFINITY;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = -INFINITY;\n for (int i = 0; i < degree; ++i) {\n int32_t src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n\n max_aggregate_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_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 5779.0s turn · 11.0s LLM · 768.0s tools · tokens in=5,070 · out=641 · total=6,054
reasoning
Still error. The error is "expected scalar type Int but found Float"
I was troubleshooting an error in the CUDA kernel implementation for the PyTorch model. The error said "expected scalar type Int but found Float," which pointed to a type mismatch.
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#include <cstdint>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int32_t start = row_ptr[dst];\n int32_t end = row_ptr[dst + 1];\n int degree = end - start;\n\n float neg_inf = -std::numeric_limits<float>::infinity();\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = neg_inf;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = neg_inf;\n for (int i = 0; i < degree; ++i) {\n int32_t src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n TORCH_CHECK(row_ptr.is_cuda() && col_idx.is_cuda() && node_feat.is_cuda(), \"Inputs must be CUDA\");\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n\n max_aggregate_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_feat)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 6209.7s turn · 5.5s LLM · 204.2s tools · tokens in=6,071 · out=642 · total=6,717
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\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cstdint>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int32_t start = row_ptr[dst];\n int32_t end = row_ptr[dst + 1];\n int degree = end - start;\n\n float neg_inf = -std::numeric_limits<float>::infinity();\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = neg_inf;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = neg_inf;\n for (int i = 0; i < degree; ++i) {\n int32_t src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n TORCH_CHECK(row_ptr.is_cuda() && col_idx.is_cuda() && node_feat.is_cuda(), \"Inputs must be CUDA\");\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n\n max_aggregate_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_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 7687.3s turn · 19.7s LLM · 667.6s tools · tokens in=6,746 · out=641 · total=7,744 · FINAL
reasoning
Still failing. Perhaps the inputs in test are on CPU?
I was troubleshooting why my custom CUDA kernel implementation for the PyTorch model wasn't working. The error indicated a scalar type mismatch, specifically "expected scalar type Int but in found Float". I suspected the tensors like row_ptr, which should be int32, might be treated as float somehow.
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#include <cstdint>\n#include <limits>\n\n__global__ void max_aggregate_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int32_t start = row_ptr[dst];\n int32_t end = row_ptr[dst + 1];\n int degree = end - start;\n\n float neg_inf = -std::numeric_limits<float>::infinity();\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = neg_inf;\n }\n return;\n }\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float max_val = neg_inf;\n for (int i = 0; i < degree; ++i) {\n int32_t src = col_idx[start + i];\n float val = node_feat[src * feat_dim + f];\n if (val > max_val) max_val = val;\n }\n out[dst * feat_dim + f] = max_val;\n }\n}\n\ntorch::Tensor max_aggregate_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor node_feat\n) {\n TORCH_CHECK(row_ptr.is_cuda() && col_idx.is_cuda() && node_feat.is_cuda(), \"Inputs must be CUDA\");\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n\n auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());\n\n int threads = 256;\n dim3 blocks(num_nodes);\n\n max_aggregate_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name='max_aggregate',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['max_aggregate_cuda'],\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, row_ptr, col_idx, node_feat):\n return self.module.max_aggregate_cuda(row_ptr, col_idx, node_feat)\n"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: correctness check did not pass (? trials).
runtime_error: expected scalar type Int but found Float
final submitted kernel (93 lines, 2,403 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>
#include <cstdint>
#include <limits>
__global__ void max_aggregate_kernel(
const int32_t* __restrict__ row_ptr,
const int32_t* __restrict__ col_idx,
const float* __restrict__ node_feat,
float* __restrict__ out,
int num_nodes,
int feat_dim
) {
int dst = blockIdx.x;
if (dst >= num_nodes) return;
int32_t start = row_ptr[dst];
int32_t end = row_ptr[dst + 1];
int degree = end - start;
float neg_inf = -std::numeric_limits<float>::infinity();
if (degree == 0) {
for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {
out[dst * feat_dim + f] = neg_inf;
}
return;
}
for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {
float max_val = neg_inf;
for (int i = 0; i < degree; ++i) {
int32_t src = col_idx[start + i];
float val = node_feat[src * feat_dim + f];
if (val > max_val) max_val = val;
}
out[dst * feat_dim + f] = max_val;
}
}
torch::Tensor max_aggregate_cuda(
torch::Tensor row_ptr,
torch::Tensor col_idx,
torch::Tensor node_feat
) {
TORCH_CHECK(row_ptr.is_cuda() && col_idx.is_cuda() && node_feat.is_cuda(), "Inputs must be CUDA");
int num_nodes = row_ptr.size(0) - 1;
int feat_dim = node_feat.size(1);
auto out = torch::empty({num_nodes, feat_dim}, node_feat.options());
int threads = 256;
dim3 blocks(num_nodes);
max_aggregate_kernel<<<blocks, threads>>>(
row_ptr.data_ptr<int32_t>(),
col_idx.data_ptr<int32_t>(),
node_feat.data_ptr<float>(),
out.data_ptr<float>(),
num_nodes,
feat_dim
);
return out;
}
"""
cpp_source = """
torch::Tensor max_aggregate_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor node_feat);
"""
module = load_inline(
name='max_aggregate',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['max_aggregate_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
self.module = module
def forward(self, row_ptr, col_idx, node_feat):
return self.module.max_aggregate_cuda(row_ptr, col_idx, node_feat)