grok-4.3 · popcorn · pop_l123_default_gpt
incorrect 18_CSRMultiHeadSpMM.py
model
grok-4.3
turns / max
8 / 12
tool calls
8
tokens (total)
41824
wall clock
1464.2s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
25.90 μs
speedup
—
started
2026-05-11T11:48:44Z
correctness issue:
timingkernelbench.timing
kernel runtime
—
ref runtime
25.90 μs
speedup
—
kernel mean / median / std
—
ref mean / median / std
25.90 / — / 4.97 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
4.000e-04
dram util %
0.0%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
1.47 GB/s
dram util %
0.0%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
6.2%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
101.0%
ld sectors / req
0.00
st sectors / req
—
regs / thread
16
smem / block
1024 B
block size
128
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
short_scoreboard
385.7%
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Multi-head sparse aggregation with per-edge attention weights.

    Each edge carries one weight per head, and source node features are arranged
    as `(num_nodes, num_heads, head_dim)`. The output aggregates weighted
    source features into `(num_nodes, num_heads, head_dim)`.
    """

    def __init__(self):
        super().__init__()

    def forward(
        self,
        row_ptr: torch.Tensor,
        col_idx: torch.Tensor,
        edge_weight: torch.Tensor,
        node_feat: torch.Tensor,
    ) -> torch.Tensor:
        num_nodes = row_ptr.numel() - 1
        num_heads = node_feat.shape[1]
        head_dim = node_feat.shape[2]
        out = torch.zeros(num_nodes, num_heads, head_dim, 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:
                src = col_idx[start:end].long()
                weights = edge_weight[start:end].unsqueeze(-1)
                out[dst] = (node_feat[src] * weights).sum(dim=0)

        return out


num_nodes = 320
avg_degree = 14
num_heads = 8
head_dim = 32


def get_inputs():
    return list(
        popcorn_pri.make_csr_multihead_spmm_graph(
            num_nodes, avg_degree, num_heads, head_dim
        )
    )



def get_init_inputs():
    return []
turn 0431.9s turn · 5.5s LLM · 426.3s tools · tokens in=1,712 · out=6 · total=2,218
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model does multi-head sparse aggregation with per-edge attention weights. The task involves replacing the forward pass of a PyTorch model with a custom CUDA kernel. The model performs multi-head sparse aggregation with per-edge attention weights, using a CSR-like format for the graph: row_ptr and col_idx define edges, edge_weight is per edge per head, and node_feat is shaped appropriately.
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 1239.4s turn · 20.9s LLM · 218.5s tools · tokens in=2,451 · out=639 · total=3,412
reasoning
GPU is H100, high bandwidth, but graph is tiny so probably compute bound or launch overhead, but anyway.
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 sparse_agg_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim,\n int num_edges\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + h * head_dim + d] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n int num_edges = col_idx.numel();\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim; // 256\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim,\n num_edges\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2234.9s turn · 6.3s LLM · 228.6s tools · tokens in=3,429 · out=640 · total=4,076
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\n__global__ void sparse_agg_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim,\n int num_edges\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + h * head_dim + d] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n int num_edges = col_idx.numel();\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim; // 256\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim,\n num_edges\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, 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 391.6s turn · 8.2s LLM · 83.4s tools · tokens in=4,105 · out=625 · total=4,934
reasoning
Error: row_ptr is int32, but in kernel I used int which on 64-bit is 64-bit? No, data_ptr<int>() expects int32? In CUDA, int is 32-bit.
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\n__global__ void sparse_agg_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int32_t src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + tid] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim;\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 433.4s turn · 8.9s LLM · 24.6s tools · tokens in=4,951 · out=626 · total=5,581
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\n__global__ void sparse_agg_kernel(\n const int32_t* __restrict__ row_ptr,\n const int32_t* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int32_t src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + tid] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim;\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n col_idx.data_ptr<int32_t>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, 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 5178.5s turn · 13.3s LLM · 165.2s tools · tokens in=5,610 · out=610 · total=6,464
reasoning
Still error. The error "expected scalar type Int but found Float" suggests perhaps row_ptr is being treated wrong, or maybe in the reference? 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\n__global__ void sparse_agg_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + tid] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim;\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 688.7s turn · 5.5s LLM · 83.3s tools · tokens in=6,481 · out=611 · total=7,096
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\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void sparse_agg_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + tid] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim;\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, 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 7165.8s turn · 13.0s LLM · 152.8s tools · tokens in=7,125 · out=610 · total=8,043 · FINAL
reasoning
Still same. The error might be coming from the reference implementation because in reference, row_ptr[dst].item() , if row_ptr is int32 fine. I noticed the error message about expecting an Int but finding a Float, which suggests a type mismatch somewhere. In the reference implementation, row_ptr is defined as int32, and operations like row_ptr[dst].
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 sparse_agg_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_weight,\n const float* __restrict__ node_feat,\n float* __restrict__ out,\n int num_nodes,\n int num_heads,\n int head_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int tid = threadIdx.x;\n int total_dim = num_heads * head_dim;\n if (tid >= total_dim) return;\n\n int h = tid / head_dim;\n int d = tid % head_dim;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n\n float sum = 0.0f;\n for (int e = start; e < end; ++e) {\n int src = col_idx[e];\n float w = edge_weight[e * num_heads + h];\n float feat = node_feat[src * total_dim + h * head_dim + d];\n sum += w * feat;\n }\n\n out[dst * total_dim + tid] = sum;\n}\n\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n int num_nodes = row_ptr.numel() - 1;\n int num_heads = node_feat.size(1);\n int head_dim = node_feat.size(2);\n\n auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());\n\n int threads = num_heads * head_dim;\n dim3 blocks(num_nodes);\n\n sparse_agg_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_weight.data_ptr<float>(),\n node_feat.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n num_heads,\n head_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_agg_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_agg\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_agg_cuda\"],\n verbose=True,\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, edge_weight, node_feat):\n return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, 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 (98 lines, 2,337 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 sparse_agg_kernel(
    const int* __restrict__ row_ptr,
    const int* __restrict__ col_idx,
    const float* __restrict__ edge_weight,
    const float* __restrict__ node_feat,
    float* __restrict__ out,
    int num_nodes,
    int num_heads,
    int head_dim
) {
    int dst = blockIdx.x;
    if (dst >= num_nodes) return;

    int tid = threadIdx.x;
    int total_dim = num_heads * head_dim;
    if (tid >= total_dim) return;

    int h = tid / head_dim;
    int d = tid % head_dim;

    int start = row_ptr[dst];
    int end = row_ptr[dst + 1];

    float sum = 0.0f;
    for (int e = start; e < end; ++e) {
        int src = col_idx[e];
        float w = edge_weight[e * num_heads + h];
        float feat = node_feat[src * total_dim + h * head_dim + d];
        sum += w * feat;
    }

    out[dst * total_dim + tid] = sum;
}

torch::Tensor sparse_agg_cuda(
    torch::Tensor row_ptr,
    torch::Tensor col_idx,
    torch::Tensor edge_weight,
    torch::Tensor node_feat
) {
    int num_nodes = row_ptr.numel() - 1;
    int num_heads = node_feat.size(1);
    int head_dim = node_feat.size(2);

    auto out = torch::zeros({num_nodes, num_heads, head_dim}, node_feat.options());

    int threads = num_heads * head_dim;
    dim3 blocks(num_nodes);

    sparse_agg_kernel<<<blocks, threads>>>(
        row_ptr.data_ptr<int>(),
        col_idx.data_ptr<int>(),
        edge_weight.data_ptr<float>(),
        node_feat.data_ptr<float>(),
        out.data_ptr<float>(),
        num_nodes,
        num_heads,
        head_dim
    );

    return out;
}
"""

cpp_source = """
torch::Tensor sparse_agg_cuda(
    torch::Tensor row_ptr,
    torch::Tensor col_idx,
    torch::Tensor edge_weight,
    torch::Tensor node_feat
);
"""

module = load_inline(
    name="sparse_agg",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["sparse_agg_cuda"],
    verbose=True,
)

class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()
        self.module = module

    def forward(self, row_ptr, col_idx, edge_weight, node_feat):
        return self.module.sparse_agg_cuda(row_ptr, col_idx, edge_weight, node_feat)