grok-4.3 · popcorn · pop_l123_default_gpt
correct 3_CSRSpMMMessagePassing.py
model
grok-4.3
turns / max
7 / 12
tool calls
7
tokens (total)
29010
wall clock
5077.4s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.04 μs
ref runtime
1860.00 μs
speedup
41425.39x
started
2026-05-11T14:04:04Z
timingkernelbench.timing
kernel runtime
0.04 μs
ref runtime
1860.00 μs
speedup
41425.39x
kernel mean / median / std
0.04 / — / 0.24 μs
ref mean / median / std
1860.00 / — / 0.05 μs
numerical precisioneval_kernel_against_ref
max abs error
4.77e-06
mean abs error
3.15e-07
max rel error
0.02
mean rel error
5.38e-07
gpu memoryextended_metrics.measure_memory
kernel peak
33.75 MB
ref peak
33.70 MB
kernel / ref
1.00x
kernel launches / fusiontorch.profiler
kernel launches
12
ref launches
21508
fusion ratio (ref / kernel)
1792.33x
total cuda time
7365 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.081
dram util %
8.1%
compute util %
7.8%
bottleneck
memory-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
270.85 GB/s
dram util %
8.1%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
48.0%
bottleneck
memory-bound
dominant pipe %
7.8%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
377.4%
ld sectors / req
9.85
st sectors / req
—
regs / thread
112
smem / block
4096 B
block size
640
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
4535.0%
short_scoreboard
664.0%
mio_throttle
149.4%
flagssubmit_kernel guard
excessive speedup
flagged
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Sparse message passing with CSR adjacency.

    For each destination node, aggregate weighted source-node features from the
    incoming edges in its CSR row. This is a core primitive in sparse graph
    neural network layers and graph convolution operators.
    """

    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
        feat_dim = node_feat.shape[1]
        out = torch.zeros(num_nodes, feat_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 = 768
avg_degree = 20
feat_dim = 96




def get_inputs():
    return list(popcorn_pri.make_csr_graph_with_edge_weights(num_nodes, avg_degree, feat_dim))



def get_init_inputs():
    return []
turn 0356.2s turn · 8.8s LLM · 347.4s tools · tokens in=1,686 · out=6 · total=1,950
reasoning
The task is to replace the forward computation with a custom CUDA kernel for sparse message passing with CSR adjacency.
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 1775.4s turn · 6.3s LLM · 769.1s tools · tokens in=2,183 · out=557 · total=2,820
reasoning
GPU is H100, high bandwidth.
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 spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_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.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256); // or min(256, feat_dim)\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n return self.spmm(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2949.9s turn · 4.9s LLM · 945.0s tools · tokens in=2,837 · out=558 · total=3,402
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 spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_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.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256); // or min(256, feat_dim)\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n return self.spmm(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 3581.1s turn · 6.7s LLM · 574.4s tools · tokens in=3,431 · out=562 · total=4,155
reasoning
Error: expected scalar type Int but found Float. Probably row_ptr or col_idx dtype.
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 spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_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.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256);\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n return self.spmm(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 4858.1s turn · 5.8s LLM · 852.3s tools · tokens in=4,172 · out=563 · total=4,740
reasoning
Now check 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#include <cstdint>\n\n__global__ void spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_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.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256);\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n return self.spmm(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 5585.4s turn · 13.9s LLM · 571.6s tools · tokens in=4,769 · out=637 · total=5,607
reasoning
Still the same error. Hmm.
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 spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n TORCH_CHECK(row_ptr.dtype() == torch::kInt32, \"row_ptr must be int32\");\n TORCH_CHECK(col_idx.dtype() == torch::kInt32, \"col_idx must be int32\");\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256);\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n col_idx = col_idx.to(torch.int32).contiguous()\n return self.spmm(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=4.77e-06 mean_abs_err=3.15e-07 max_rel_err=2.21e-02
turn 6971.3s turn · 7.3s LLM · 964.0s tools · tokens in=5,667 · out=636 · total=6,336 · FINAL
reasoning
Great, now it's correct. The casts were needed perhaps due to some dtype issue in tests.
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\n__global__ void spmm_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 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\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\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];\n sum += node_feat[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = sum;\n }\n}\n\ntorch::Tensor spmm_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_weight,\n torch::Tensor node_feat\n) {\n TORCH_CHECK(row_ptr.dtype() == torch::kInt32, \"row_ptr must be int32\");\n TORCH_CHECK(col_idx.dtype() == torch::kInt32, \"col_idx must be int32\");\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_feat.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(256);\n\n spmm_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 feat_dim\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor spmm_cuda(torch::Tensor row_ptr, torch::Tensor col_idx, torch::Tensor edge_weight, torch::Tensor node_feat);\n\"\"\"\n\nmodule = load_inline(\n name=\"spmm_module\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"spmm_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.spmm = module.spmm_cuda\n\n def forward(self, row_ptr, col_idx, edge_weight, node_feat):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n col_idx = col_idx.to(torch.int32).contiguous()\n return self.spmm(row_ptr, col_idx, edge_weight, node_feat)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.04 μs Runtime stats: mean=0.04μs median=?μs std=0.24μs Numerical precision: max_abs_err=4.77e-06 mean_abs_err=3.15e-07 max_rel_err=2.21e-02 Memory: 33.75 MB peak (ref: 33.7 MB, ratio: 1.0013x) Kernel launches: 12 (ref: 21508, fusion ratio: 1792.3333) SOL score: 0.0809 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (87 lines, 2,348 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>

__global__ void spmm_kernel(
    const int32_t* __restrict__ row_ptr,
    const int32_t* __restrict__ col_idx,
    const float* __restrict__ edge_weight,
    const float* __restrict__ node_feat,
    float* __restrict__ out,
    int num_nodes,
    int feat_dim
) {
    int dst = blockIdx.x;
    if (dst >= num_nodes) return;

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

    for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {
        float sum = 0.0f;
        for (int e = start; e < end; ++e) {
            int32_t src = col_idx[e];
            float w = edge_weight[e];
            sum += node_feat[src * feat_dim + f] * w;
        }
        out[dst * feat_dim + f] = sum;
    }
}

torch::Tensor spmm_cuda(
    torch::Tensor row_ptr,
    torch::Tensor col_idx,
    torch::Tensor edge_weight,
    torch::Tensor node_feat
) {
    TORCH_CHECK(row_ptr.dtype() == torch::kInt32, "row_ptr must be int32");
    TORCH_CHECK(col_idx.dtype() == torch::kInt32, "col_idx must be int32");
    int num_nodes = row_ptr.size(0) - 1;
    int feat_dim = node_feat.size(1);
    auto out = torch::zeros({num_nodes, feat_dim}, node_feat.options());

    dim3 blocks(num_nodes);
    dim3 threads(256);

    spmm_kernel<<<blocks, threads>>>(
        row_ptr.data_ptr<int32_t>(),
        col_idx.data_ptr<int32_t>(),
        edge_weight.data_ptr<float>(),
        node_feat.data_ptr<float>(),
        out.data_ptr<float>(),
        num_nodes,
        feat_dim
    );

    return out;
}
"""

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

module = load_inline(
    name="spmm_module",
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=["spmm_cuda"],
    verbose=False
)

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

    def forward(self, row_ptr, col_idx, edge_weight, node_feat):
        row_ptr = row_ptr.to(torch.int32).contiguous()
        col_idx = col_idx.to(torch.int32).contiguous()
        return self.spmm(row_ptr, col_idx, edge_weight, node_feat)