grok-4.3 · popcorn · pop_l123_default_gpt
correct 8_SegmentTopKCSR.py
model
grok-4.3
turns / max
12 / 12
tool calls
12
tokens (total)
97471
wall clock
5208.6s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.04 μs
ref runtime
22.10 μs
speedup
602.18x
started
2026-05-11T13:50:17Z
timingkernelbench.timing
kernel runtime
0.04 μs
ref runtime
22.10 μs
speedup
602.18x
kernel mean / median / std
0.04 / — / 0.00 μs
ref mean / median / std
22.10 / — / 0.73 μs
numerical precisioneval_kernel_against_ref
max abs error
0.00
mean abs error
0.00
max rel error
0.00
mean rel error
0.00
gpu memoryextended_metrics.measure_memory
kernel peak
33.03 MB
ref peak
33.03 MB
kernel / ref
1.00x
kernel launches / fusiontorch.profiler
kernel launches
6
ref launches
6086
fusion ratio (ref / kernel)
1014.33x
total cuda time
251 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.003
dram util %
0.3%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
9.27 GB/s
dram util %
0.3%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
14.0%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
194.9%
ld sectors / req
26.82
st sectors / req
—
regs / thread
72
smem / block
2048 B
block size
384
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
1745.5%
short_scoreboard
97.5%
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):
    """
    Per-row top-k selection over CSR edge segments.

    Returns the top-k values and their original edge indices for each node's
    neighborhood. Useful for graph sparsification, hard attention, and
    candidate-pruning workloads.
    """

    def __init__(self, k):
        super().__init__()
        self.k = k

    def forward(
        self,
        row_ptr: torch.Tensor,
        edge_scores: torch.Tensor,
    ) -> tuple[torch.Tensor, torch.Tensor]:
        num_nodes = row_ptr.numel() - 1
        topk_vals = torch.full((num_nodes, self.k), float("-inf"), dtype=edge_scores.dtype, device=edge_scores.device)
        topk_idx = torch.full((num_nodes, self.k), -1, dtype=torch.int64, device=edge_scores.device)

        for node in range(num_nodes):
            start = int(row_ptr[node].item())
            end = int(row_ptr[node + 1].item())
            if end <= start:
                continue
            segment = edge_scores[start:end]
            curr_k = min(self.k, end - start)
            vals, idx = torch.topk(segment, k=curr_k, dim=0)
            topk_vals[node, :curr_k] = vals
            topk_idx[node, :curr_k] = idx.long() + start

        return topk_vals, topk_idx


num_nodes = 320
avg_degree = 12
k = 4




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



def get_init_inputs():
    return [k]
turn 01608.8s turn · 10.0s LLM · 1598.8s tools · tokens in=1,707 · out=6 · total=2,141
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model does per-row top-k selection over CSR edge segments. I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs per-row top-k selection over CSR edge segments. The reference implementation loops over each node, sorts or uses torch.topk on each node's edge_scores segment, and outputs the concatenated top-k values and indices.
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 1314.7s turn · 18.4s LLM · 296.3s tools · tokens in=2,374 · out=936 · total=3,505
reasoning
GPU is H100, high bandwidth, but data is tiny: row_ptr 321 ints ~1KB, scores 4k floats ~16KB, outputs 320*4*2 ~5KB. So memory bound trivially, but reference is slow due to Python loop.
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 <vector>\n#include <limits>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n // Simple selection for top-k using one thread\n // We will find top min(k, degree)\n int curr_k = min(k, degree);\n\n // Use local arrays for top values and indices\n float top_vals[16]; // assume k <=16\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n // insert into top if larger\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n // shift\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=True,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2223.4s turn · 5.9s LLM · 217.4s tools · tokens in=3,522 · out=937 · total=4,466
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 <vector>\n#include <limits>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n // Simple selection for top-k using one thread\n // We will find top min(k, degree)\n int curr_k = min(k, degree);\n\n // Use local arrays for top values and indices\n float top_vals[16]; // assume k <=16\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n // insert into top if larger\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n // shift\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\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 3572.9s turn · 11.2s LLM · 561.6s tools · tokens in=4,495 · out=904 · total=5,540
reasoning
Error: row_ptr is int32 but kernel expects int which is 32-bit ok, but wait, error "expected scalar type Int but found Float" – probably in data_ptr, row_ptr is int32, but in C++ I used int*, torch in...
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 <vector>\n#include <limits>\n#include <cstdint>\n\n__global__ void topk_kernel(\n const int32_t* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int32_t start = row_ptr[node];\n int32_t end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 4194.2s turn · 8.1s LLM · 186.1s tools · tokens in=5,557 · out=905 · total=6,466
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 <vector>\n#include <limits>\n#include <cstdint>\n\n__global__ void topk_kernel(\n const int32_t* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int32_t start = row_ptr[node];\n int32_t end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int32_t>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\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 5359.8s turn · 9.9s LLM · 349.9s tools · tokens in=6,495 · out=910 · total=7,768
reasoning
Still error. The error is probably inside, perhaps row_ptr is torch.int32, but in forward call ok. I was having trouble with a PyTorch model where the forward computation needed to be replaced with a custom CUDA kernel. The task was to create a `ModelNew` class that mimics the original `Model`, accepting the same inputs and signature, and producing the same outputs numerically.
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 <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n TORCH_CHECK(row_ptr.dtype() == torch::kInt32, \"row_ptr must be int32\");\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 6175.9s turn · 7.6s LLM · 168.3s tools · tokens in=7,785 · out=911 · total=8,700
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 <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n TORCH_CHECK(row_ptr.dtype() == torch::kInt32, \"row_ptr must be int32\");\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← run_correctness (fail)
failrun_correctness run_correctness FAILED: ? trials did not all match. runtime_error: row_ptr must be int32
turn 7250.9s turn · 10.9s LLM · 240.0s tools · tokens in=8,728 · out=887 · total=9,737
reasoning
Ah, the TORCH_CHECK fails, meaning row_ptr.dtype() != torch::kInt32.
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 <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 8397.6s turn · 9.6s LLM · 388.0s tools · tokens in=9,754 · out=888 · total=10,645
reasoning
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#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n return self.topk_cuda(row_ptr, edge_scores, self.k)\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 9327.4s turn · 9.7s LLM · 317.7s tools · tokens in=10,708 · out=901 · total=11,844
reasoning
Still error. The error "expected scalar type Int but found Float" suggests somewhere data_ptr<int>() on a float tensor.
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 <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 10169.0s turn · 11.3s LLM · 157.7s tools · tokens in=11,895 · out=902 · total=12,801
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#include <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n return self.topk_cuda(row_ptr, edge_scores, self.k)\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 11614.0s turn · 8.7s LLM · 605.2s tools · tokens in=12,951 · out=901 · total=13,858 · FINAL
reasoning
Now it's correct, submit.
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 <vector>\n\n__global__ void topk_kernel(\n const int* __restrict__ row_ptr,\n const float* __restrict__ edge_scores,\n float* __restrict__ topk_vals,\n int64_t* __restrict__ topk_idx,\n int num_nodes,\n int k\n) {\n int node = blockIdx.x * blockDim.x + threadIdx.x;\n if (node >= num_nodes) return;\n\n int start = row_ptr[node];\n int end = row_ptr[node + 1];\n int degree = end - start;\n\n if (degree <= 0) {\n for (int i = 0; i < k; ++i) {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n return;\n }\n\n int curr_k = min(k, degree);\n\n float top_vals[16];\n int64_t top_indices[16];\n for (int i = 0; i < curr_k; ++i) {\n top_vals[i] = -INFINITY;\n top_indices[i] = -1;\n }\n\n for (int i = 0; i < degree; ++i) {\n float val = edge_scores[start + i];\n for (int j = 0; j < curr_k; ++j) {\n if (val > top_vals[j]) {\n for (int m = curr_k - 1; m > j; --m) {\n top_vals[m] = top_vals[m-1];\n top_indices[m] = top_indices[m-1];\n }\n top_vals[j] = val;\n top_indices[j] = start + i;\n break;\n }\n }\n }\n\n for (int i = 0; i < k; ++i) {\n if (i < curr_k) {\n topk_vals[node * k + i] = top_vals[i];\n topk_idx[node * k + i] = top_indices[i];\n } else {\n topk_vals[node * k + i] = -INFINITY;\n topk_idx[node * k + i] = -1;\n }\n }\n}\n\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {\n int num_nodes = row_ptr.size(0) - 1;\n auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());\n auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());\n auto topk_vals = torch::empty({num_nodes, k}, options_val);\n auto topk_idx = torch::empty({num_nodes, k}, options_idx);\n\n int threads = 256;\n int blocks = (num_nodes + threads - 1) / threads;\n topk_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n topk_vals.data_ptr<float>(),\n topk_idx.data_ptr<int64_t>(),\n num_nodes,\n k\n );\n cudaDeviceSynchronize();\n return {topk_vals, topk_idx};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);\n'''\n\nmodule = load_inline(\n name='topk_module',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['topk_cuda'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, k):\n super().__init__()\n self.k = k\n self.topk_cuda = module.topk_cuda\n\n def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n return self.topk_cuda(row_ptr, edge_scores, self.k)\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.00μs Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00 Memory: 33.03 MB peak (ref: 33.03 MB, ratio: 1.0x) Kernel launches: 6 (ref: 6086, fusion ratio: 1014.3333) SOL score: 0.0028 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (114 lines, 3,232 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 <vector>

__global__ void topk_kernel(
    const int* __restrict__ row_ptr,
    const float* __restrict__ edge_scores,
    float* __restrict__ topk_vals,
    int64_t* __restrict__ topk_idx,
    int num_nodes,
    int k
) {
    int node = blockIdx.x * blockDim.x + threadIdx.x;
    if (node >= num_nodes) return;

    int start = row_ptr[node];
    int end = row_ptr[node + 1];
    int degree = end - start;

    if (degree <= 0) {
        for (int i = 0; i < k; ++i) {
            topk_vals[node * k + i] = -INFINITY;
            topk_idx[node * k + i] = -1;
        }
        return;
    }

    int curr_k = min(k, degree);

    float top_vals[16];
    int64_t top_indices[16];
    for (int i = 0; i < curr_k; ++i) {
        top_vals[i] = -INFINITY;
        top_indices[i] = -1;
    }

    for (int i = 0; i < degree; ++i) {
        float val = edge_scores[start + i];
        for (int j = 0; j < curr_k; ++j) {
            if (val > top_vals[j]) {
                for (int m = curr_k - 1; m > j; --m) {
                    top_vals[m] = top_vals[m-1];
                    top_indices[m] = top_indices[m-1];
                }
                top_vals[j] = val;
                top_indices[j] = start + i;
                break;
            }
        }
    }

    for (int i = 0; i < k; ++i) {
        if (i < curr_k) {
            topk_vals[node * k + i] = top_vals[i];
            topk_idx[node * k + i] = top_indices[i];
        } else {
            topk_vals[node * k + i] = -INFINITY;
            topk_idx[node * k + i] = -1;
        }
    }
}

std::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k) {
    int num_nodes = row_ptr.size(0) - 1;
    auto options_val = torch::TensorOptions().dtype(edge_scores.dtype()).device(edge_scores.device());
    auto options_idx = torch::TensorOptions().dtype(torch::kInt64).device(edge_scores.device());
    auto topk_vals = torch::empty({num_nodes, k}, options_val);
    auto topk_idx = torch::empty({num_nodes, k}, options_idx);

    int threads = 256;
    int blocks = (num_nodes + threads - 1) / threads;
    topk_kernel<<<blocks, threads>>>(
        row_ptr.data_ptr<int>(),
        edge_scores.data_ptr<float>(),
        topk_vals.data_ptr<float>(),
        topk_idx.data_ptr<int64_t>(),
        num_nodes,
        k
    );
    cudaDeviceSynchronize();
    return {topk_vals, topk_idx};
}
'''

cpp_source = '''
std::vector<torch::Tensor> topk_cuda(torch::Tensor row_ptr, torch::Tensor edge_scores, int k);
'''

module = load_inline(
    name='topk_module',
    cpp_sources=[cpp_source],
    cuda_sources=[cuda_source],
    functions=['topk_cuda'],
    verbose=False,
    extra_cflags=['-O3'],
    extra_cuda_cflags=['-O3']
)

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

    def forward(self, row_ptr: torch.Tensor, edge_scores: torch.Tensor):
        row_ptr = row_ptr.to(torch.int32).contiguous()
        return self.topk_cuda(row_ptr, edge_scores, self.k)