correct
8_SegmentTopKCSR.py
started
2026-05-11T13:50:17Z
timingkernelbench.timing
kernel mean / median / std
0.04 / — / 0.00 μs
ref mean / median / std
22.10 / — / 0.73 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
1014.33x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
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)