correct
37_CSRFusedAttentionValue.py
started
2026-05-11T18:42:34Z
timingkernelbench.timing
kernel mean / median / std
0.02 / — / 7.48e-04 μs
ref mean / median / std
28.50 / — / 0.25 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
960.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):
"""
Fused sparse attention-value aggregation in CSR format.
For each destination row, compute a softmax over edge logits and then use
the normalized weights to aggregate source node values. This is a common
fused graph-attention pattern.
"""
def __init__(self):
super().__init__()
def forward(
self,
row_ptr: torch.Tensor,
col_idx: torch.Tensor,
edge_scores: torch.Tensor,
node_value: torch.Tensor,
) -> torch.Tensor:
num_nodes = row_ptr.numel() - 1
feat_dim = node_value.shape[1]
out = torch.zeros(num_nodes, feat_dim, dtype=node_value.dtype, device=node_value.device)
for dst in range(num_nodes):
start = int(row_ptr[dst].item())
end = int(row_ptr[dst + 1].item())
if end > start:
weights = torch.softmax(edge_scores[start:end], dim=0)
src = col_idx[start:end].long()
out[dst] = (node_value[src] * weights.unsqueeze(-1)).sum(dim=0)
return out
num_nodes = 384
avg_degree = 18
feat_dim = 64
def get_inputs():
return list(
popcorn_pri.make_csr_fused_attention_value_graph(
num_nodes, avg_degree, feat_dim
)
)
def get_init_inputs():
return []
turn 0227.3s turn · 7.2s LLM · 220.1s tools · tokens in=2,211 · out=6 · total=2,611
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs fused sparse attention-value aggregation in CSR format.
The task involves replacing the forward computation of a PyTorch model with a custom CUDA kernel. The model performs a fused sparse attention-value aggregation in CSR format. The inputs are CSR pointers row_ptr, column indices col_idx, edge scores edge_scores, and node values node_value.
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 1307.3s turn · 12.3s LLM · 295.0s tools · tokens in=2,844 · out=1,148 · total=4,144
reasoning
GPU is H100, high bandwidth, good for this memory-bound op probably.
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 <cmath>\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) return;\n\n // Compute max for softmax\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n if (s > max_score) max_score = s;\n }\n // Reduce max across threads\n __shared__ float shared_max[32];\n int lane = threadIdx.x % 32;\n int warp = threadIdx.x / 32;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < (blockDim.x + 31)/32) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n // Compute sum exp\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n // Reduce sum\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < (blockDim.x + 31)/32) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n // Now aggregate\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\n__device__ float warpReduceMax(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ float warpReduceSum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64); // or 128\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.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 sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'sparse_attn_value'
turn 2232.2s turn · 15.8s LLM · 216.5s tools · tokens in=4,169 · out=1,256 · total=5,492
reasoning
Compilation failed. Probably missing warp reduce functions or something. Need to fix the code.
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 <cmath>\n\n__device__ __forceinline__ float warpReduceMax(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warpReduceSum(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = 0.0f;\n }\n return;\n }\n\n // Compute max for softmax\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n max_score = fmaxf(max_score, s);\n }\n // Reduce max across threads\n __shared__ float shared_max[32];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n // Compute sum exp\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = __expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n // Reduce sum\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n // Now aggregate - each thread computes its features\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = __expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64);\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n cudaError_t err = cudaGetLastError();\n if (err != cudaSuccess) {\n printf(\"CUDA error: %s\\\\n\", cudaGetErrorString(err));\n }\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 3316.9s turn · 13.5s LLM · 303.3s tools · tokens in=5,509 · out=1,257 · total=6,773
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 <cmath>\n\n__device__ __forceinline__ float warpReduceMax(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warpReduceSum(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = 0.0f;\n }\n return;\n }\n\n // Compute max for softmax\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n max_score = fmaxf(max_score, s);\n }\n // Reduce max across threads\n __shared__ float shared_max[32];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n // Compute sum exp\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = __expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n // Reduce sum\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n // Now aggregate - each thread computes its features\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = __expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64);\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.data_ptr<float>(),\n out.data_ptr<float>(),\n num_nodes,\n feat_dim\n );\n\n cudaError_t err = cudaGetLastError();\n if (err != cudaSuccess) {\n printf(\"CUDA error: %s\\\\n\", cudaGetErrorString(err));\n }\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\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 4307.6s turn · 21.2s LLM · 286.4s tools · tokens in=6,802 · out=1,229 · total=8,208
reasoning
Error: expected scalar type Int but found Float. Probably row_ptr or col_idx dtype issue? No, inputs are int32.
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 <cmath>\n\n__device__ __forceinline__ float warpReduceMax(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warpReduceSum(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = 0.0f;\n }\n return;\n }\n\n // Compute max for softmax\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n max_score = fmaxf(max_score, s);\n }\n // Reduce max across threads\n __shared__ float shared_max[32];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n // Compute sum exp\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = __expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n // Reduce sum\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n // Now aggregate - each thread computes its features\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = __expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64);\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.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 sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n col_idx = col_idx.to(torch.int32)\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\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 5296.3s turn · 12.7s LLM · 283.6s tools · tokens in=8,237 · out=1,210 · total=9,550
reasoning
Still error. The error is "expected scalar type Int but found Float" – this is likely from inside the kernel call? No.
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 <cmath>\n\n__device__ __forceinline__ float warpReduceMax(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warpReduceSum(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = 0.0f;\n }\n return;\n }\n\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n max_score = fmaxf(max_score, s);\n }\n __shared__ float shared_max[32];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = __expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = __expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64);\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.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 sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n col_idx = col_idx.to(torch.int32).contiguous()\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.84e-08 max_rel_err=5.66e-03
turn 6255.2s turn · 11.9s LLM · 243.2s tools · tokens in=9,610 · out=1,209 · total=10,838 · FINAL
reasoning
Correctness passed! Great.
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 <cmath>\n\n__device__ __forceinline__ float warpReduceMax(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warpReduceSum(float val) {\n #pragma unroll\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n__global__ void sparse_attention_value_kernel(\n const int* __restrict__ row_ptr,\n const int* __restrict__ col_idx,\n const float* __restrict__ edge_scores,\n const float* __restrict__ node_value,\n float* __restrict__ out,\n int num_nodes,\n int feat_dim\n) {\n int dst = blockIdx.x;\n if (dst >= num_nodes) return;\n\n int start = row_ptr[dst];\n int end = row_ptr[dst + 1];\n int degree = end - start;\n\n if (degree == 0) {\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n out[dst * feat_dim + f] = 0.0f;\n }\n return;\n }\n\n float max_score = -INFINITY;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = edge_scores[start + i];\n max_score = fmaxf(max_score, s);\n }\n __shared__ float shared_max[32];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[warp] = max_score;\n __syncthreads();\n if (warp == 0) {\n max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;\n max_score = warpReduceMax(max_score);\n if (lane == 0) shared_max[0] = max_score;\n }\n __syncthreads();\n max_score = shared_max[0];\n\n float sum_exp = 0.0f;\n for (int i = threadIdx.x; i < degree; i += blockDim.x) {\n float s = __expf(edge_scores[start + i] - max_score);\n sum_exp += s;\n }\n __shared__ float shared_sum[32];\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[warp] = sum_exp;\n __syncthreads();\n if (warp == 0) {\n sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;\n sum_exp = warpReduceSum(sum_exp);\n if (lane == 0) shared_sum[0] = sum_exp;\n }\n __syncthreads();\n sum_exp = shared_sum[0];\n float inv_sum = 1.0f / sum_exp;\n\n for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {\n float val = 0.0f;\n for (int i = 0; i < degree; i++) {\n float w = __expf(edge_scores[start + i] - max_score) * inv_sum;\n int src = col_idx[start + i];\n val += node_value[src * feat_dim + f] * w;\n }\n out[dst * feat_dim + f] = val;\n }\n}\n\ntorch::Tensor sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value\n) {\n int num_nodes = row_ptr.size(0) - 1;\n int feat_dim = node_value.size(1);\n auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());\n\n dim3 blocks(num_nodes);\n dim3 threads(64);\n\n sparse_attention_value_kernel<<<blocks, threads>>>(\n row_ptr.data_ptr<int>(),\n col_idx.data_ptr<int>(),\n edge_scores.data_ptr<float>(),\n node_value.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 sparse_attention_value_cuda(\n torch::Tensor row_ptr,\n torch::Tensor col_idx,\n torch::Tensor edge_scores,\n torch::Tensor node_value);\n\"\"\"\n\nmodule = load_inline(\n name=\"sparse_attn_value\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"sparse_attention_value_cuda\"],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, row_ptr, col_idx, edge_scores, node_value):\n row_ptr = row_ptr.to(torch.int32).contiguous()\n col_idx = col_idx.to(torch.int32).contiguous()\n return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.02 μs
Runtime stats: mean=0.02μs median=?μs std=0.00μs
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.84e-08 max_rel_err=5.66e-03
Memory: 33.27 MB peak (ref: 33.25 MB, ratio: 1.0005x)
Kernel launches: 12 (ref: 11524, fusion ratio: 960.3333)
SOL score: 0.0535
Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (150 lines, 4,291 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 <cmath>
__device__ __forceinline__ float warpReduceMax(float val) {
#pragma unroll
for (int offset = 16; offset > 0; offset /= 2) {
val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));
}
return val;
}
__device__ __forceinline__ float warpReduceSum(float val) {
#pragma unroll
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_down_sync(0xffffffff, val, offset);
}
return val;
}
__global__ void sparse_attention_value_kernel(
const int* __restrict__ row_ptr,
const int* __restrict__ col_idx,
const float* __restrict__ edge_scores,
const float* __restrict__ node_value,
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];
int degree = end - start;
if (degree == 0) {
for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {
out[dst * feat_dim + f] = 0.0f;
}
return;
}
float max_score = -INFINITY;
for (int i = threadIdx.x; i < degree; i += blockDim.x) {
float s = edge_scores[start + i];
max_score = fmaxf(max_score, s);
}
__shared__ float shared_max[32];
int lane = threadIdx.x & 31;
int warp = threadIdx.x >> 5;
max_score = warpReduceMax(max_score);
if (lane == 0) shared_max[warp] = max_score;
__syncthreads();
if (warp == 0) {
max_score = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_max[threadIdx.x] : -INFINITY;
max_score = warpReduceMax(max_score);
if (lane == 0) shared_max[0] = max_score;
}
__syncthreads();
max_score = shared_max[0];
float sum_exp = 0.0f;
for (int i = threadIdx.x; i < degree; i += blockDim.x) {
float s = __expf(edge_scores[start + i] - max_score);
sum_exp += s;
}
__shared__ float shared_sum[32];
sum_exp = warpReduceSum(sum_exp);
if (lane == 0) shared_sum[warp] = sum_exp;
__syncthreads();
if (warp == 0) {
sum_exp = (threadIdx.x < ((blockDim.x + 31) >> 5)) ? shared_sum[threadIdx.x] : 0.0f;
sum_exp = warpReduceSum(sum_exp);
if (lane == 0) shared_sum[0] = sum_exp;
}
__syncthreads();
sum_exp = shared_sum[0];
float inv_sum = 1.0f / sum_exp;
for (int f = threadIdx.x; f < feat_dim; f += blockDim.x) {
float val = 0.0f;
for (int i = 0; i < degree; i++) {
float w = __expf(edge_scores[start + i] - max_score) * inv_sum;
int src = col_idx[start + i];
val += node_value[src * feat_dim + f] * w;
}
out[dst * feat_dim + f] = val;
}
}
torch::Tensor sparse_attention_value_cuda(
torch::Tensor row_ptr,
torch::Tensor col_idx,
torch::Tensor edge_scores,
torch::Tensor node_value
) {
int num_nodes = row_ptr.size(0) - 1;
int feat_dim = node_value.size(1);
auto out = torch::zeros({num_nodes, feat_dim}, node_value.options());
dim3 blocks(num_nodes);
dim3 threads(64);
sparse_attention_value_kernel<<<blocks, threads>>>(
row_ptr.data_ptr<int>(),
col_idx.data_ptr<int>(),
edge_scores.data_ptr<float>(),
node_value.data_ptr<float>(),
out.data_ptr<float>(),
num_nodes,
feat_dim
);
return out;
}
"""
cpp_source = """
torch::Tensor sparse_attention_value_cuda(
torch::Tensor row_ptr,
torch::Tensor col_idx,
torch::Tensor edge_scores,
torch::Tensor node_value);
"""
module = load_inline(
name="sparse_attn_value",
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=["sparse_attention_value_cuda"],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
self.module = module
def forward(self, row_ptr, col_idx, edge_scores, node_value):
row_ptr = row_ptr.to(torch.int32).contiguous()
col_idx = col_idx.to(torch.int32).contiguous()
return self.module.sparse_attention_value_cuda(row_ptr, col_idx, edge_scores, node_value)