FW-GLM-5-1 · popcorn · pop_l123_st_gpt
incorrect 10_SampledDenseDenseMatmulEdges.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
4285
wall clock
645.8s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
0.05 μs
speedup
—
started
2026-05-11T09:41:45Z
correctness issue:
timingkernelbench.timing
kernel runtime
—
ref runtime
0.05 μs
speedup
—
kernel mean / median / std
—
ref mean / median / std
0.05 / — / 5.64e-04 μs
roofline / occupancynsight_hardware_counters
dram bandwidth
—
dram util %
—
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
—
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
—
ld sectors / req
—
st sectors / req
—
regs / thread
—
smem / block
—
block size
—
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Edge-only sampled dense-dense matmul.

    For each graph edge (src, dst), compute a dot product between source and
    destination node embeddings. This pattern appears in link prediction,
    sampled attention, and graph contrastive learning.
    """

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

    def forward(
        self,
        src_idx: torch.Tensor,
        dst_idx: torch.Tensor,
        src_feat: torch.Tensor,
        dst_feat: torch.Tensor,
    ) -> torch.Tensor:
        src = src_feat[src_idx.long()]
        dst = dst_feat[dst_idx.long()]
        return (src * dst).sum(dim=-1)


num_src_nodes = 1024
num_dst_nodes = 768
num_edges = 16384
feat_dim = 128


def get_inputs():
    p = popcorn_pri
    n_edges = p.jitter_int(num_edges)
    fd = p.jitter_int(feat_dim, align=8)
    ns = p.jitter_int(num_src_nodes)
    nd = p.jitter_int(num_dst_nodes)
    src_idx = torch.randint(0, ns, (n_edges,), dtype=torch.int32)
    dst_idx = torch.randint(0, nd, (n_edges,), dtype=torch.int32)
    src_feat = torch.randn(ns, fd, dtype=torch.float32)
    dst_feat = torch.randn(nd, fd, dtype=torch.float32)
    return [src_idx, dst_idx, src_feat, dst_feat]





def get_init_inputs():
    return []
turn 0 · FINAL · 1 tool call(s) · 645.8s turn · 38.6s LLM · tokens in=1,104 · out=3,181 · total=4,285 · 15,620 chars (click to expand)
turn 0645.8s turn · 38.6s LLM · 607.2s tools · tokens in=1,104 · out=3,181 · total=4,285 · FINAL
reasoning
Let me analyze this task. I need to write a custom CUDA kernel for an edge-only sampled dense-dense matmul operation. The operation: 1. For each edge (src_idx[i], dst_idx[i]), gather source and destination node embeddings 2. Compute element-wise product of the two embeddings 3. Sum across the feature dimension The reference implementation: ```python src = src_feat[src_idx.long()] dst = dst_feat[dst_idx.long()] return (src * dst).sum(dim=-1) ``` This is essentially: - src_feat is [num_src_nodes, feat_dim] - dst_feat is [num_dst_nodes, feat_dim] - src_idx is [num_edges] - indices into src_feat - dst_idx is [num_edges] - indices into dst_feat - Output is [num_edges] - each element is a dot product So for each edge i: - output[i] = sum_j(src_feat[src_idx[i], j] * dst_feat[dst_idx[i], j]) The default parameters: - num_src_nodes = 1024 - num_dst_nodes = 768 - num_edges = 16384 - feat_dim = 128 I need to write an efficient CUDA kernel for this. Let me think about the approach: 1. Each thread can handle one edge - this means num_edges = 16384 threads 2. For each edge, we need to compute the dot product of two feature vectors of length 128 One approach: Each thread handles one edge, iterates through the feature dimension - Total threads: num_edges - Each thread does 128 multiplications and 128 additions - This might be memory bound since we're accessing src_feat and dst_feat Another approach: Use shared memory or warp-level parallelism for the reduction Since feat_dim = 128 and warp size = 32, I could have each warp handle multiple edges (like 4 edges per warp, where each thread handles 32 features for one edge), or each block handles multiple edges with shared memory. Looking at the memory access patterns, src_idx[i] determines which row of src_feat to read, and dst_idx[i] determines which row of dst_feat to read. The feature dimension is indexed by j, so coalescing memory reads would be important for performance. With one thread per edge and sequential feature accesses, we get coalesced reads within each row of src_feat and dst_feat. But consecutive threads access different rows, which could lead to poor cache locality. Using one thread per feature per edge instead would enable better coalescing: threads in a warp read consecutive elements from the same feature row, reducing cache misses. After computing the products, I'd need a parallel reduction across the feature dimension. The simplest approach starts with one thread per edge handling the entire dot product sequentially. For optimization, I could assign multiple threads per edge to parallelize the feature dimension computation, then perform a warp or block reduction to get the final result. The key insight is that consecutive threads accessing the same feature row enables coalesced memory access, which should significantly improve performance on 16384 edges with 128 features. For a basic implementation, each thread processes one edge independently, computing the dot product by iterating through all 128 features. The kernel uses grid-stride loops to handle variable edge counts efficiently, loading source and destination features for each edge and accumulating the product directly without shared memory overhead. However, this approach suffers from memory access inefficiency—consecutive threads read from scattered locations in src_feat and dst_feat based on their edge indices, leading to poor cache locality. A better strategy reorganizes the computation so that threads in a warp access consecutive features from the same edge, enabling coalesced memory reads where thread k reads feature k from both src_feat and dst_feat for a given edge. This can be implemented by assigning a warp to each edge, with each thread handling one feature element. The warp then performs a reduction using warp shuffle operations to sum across all features. For cases where the feature dimension exceeds the warp size, each thread processes multiple features in a loop before the final reduction. In the backward pass, gradients accumulate to specific rows of the feature matrices, requiring atomic additions when multiple edges reference the same node. The implementation uses a custom CUDA kernel with `torch.utils.cpp_extension.load_inline`. A warp-level shuffle reduction efficiently sums the dot product across the warp dimension. The kernel itself takes edge indices and feature matrices as input, using grid-stride loops to distribute work across warps. Each warp handles one edge, with lane IDs and mask calculations to determine active threads. Shared memory buffers temporarily store the feature indices before they're loaded into registers for computation. For each edge, threads cooperatively load source and destination features using a grid-stride loop that processes multiple features per thread. The dot product accumulates across all features, then a warp shuffle reduction sums the partial results across lanes. Thread 0 writes the final output, and the kernel advances to the next edge in the grid-stride manner. The wrapper functions bind the kernel to PyTorch tensors, using bitwise AND for edge count to satisfy the constraint that num_edges must be a power of 2. The forward pass invokes the kernel with a block size of 32 threads per warp and a grid size covering all edges. In the ModelNew class, I call the forward function through the loaded extension. The cuda_src contains the kernel implementation with the warp reduction helper using __shfl_down_sync for efficient parallel summation across threads in a warp. The kernel itself loads source and destination indices from the input arrays, then processes features in a strided loop where each thread handles a subset of the feature dimension. The dot product is accumulated across all features before being reduced across the warp. The reduction uses warp shuffle operations to sum the partial results, and the final output is written by the first thread in each warp. The wrapper function then launches this kernel with appropriate block and grid dimensions based on the number of edges. d to have enough warps. With 16384 edges and 32 threads per block (1 warp), we'd have 16384 blocks. That's actually a reasonable number for a GPU. But 1024 threads per block (32 warps) would give 512 blocks, which might be better. Actually, looking at occupancy considerations, having 32 threads per block means we can have many blocks active simultaneously. With 16384 blocks, the GPU will have plenty of work to schedule. This should be fine. Let me reconsider the memory access pattern with this kernel: - Edge i is handled by warp i (threads 0-31) - Thread k in warp i reads src_feat[src_idx[i], k] - Thread k+1 reads src_feat[src_idx[i], k+1] - These are consecutive memory locations, so they're coalesced! This is much better than the naive approach where each thread handles one edge. Let me also think about the grid-stride loop. With 16384 edges and a grid of say 512 blocks x 1 warp = 512 warps, the loop would run 16384/512 = 32 iterations. Each iteration processes 512 edges. This is good because: 1. We can scale the grid size based on GPU resources 2. The loop adds minimal overhead compared to the computation I need to make sure the types align correctly. The indices are int32, so the kernel expects `int*` for src_idx and dst_idx. The features and output are float32, matching `float*` types. Using `src_idx[i]` and `dst_idx[i]` directly as indices should work since they're already int32. For the grid configuration, I'm launching blocks with 32 threads each (one warp per edge). The grid size should be ceil(num_edges / 1), so I need to calculate this properly when launching the kernel. Using `(num_edges + threads_per_block - 1) / threads_per_block` for the grid dimension ensures I have enough blocks to cover all edges, though the grid-stride loop handles any overflow gracefully. I should account for num_edges not being a power of 2 by using proper ceiling division for the grid calculation. For the kernel itself, I'm using a simple dot product approach where each thread processes one edge. This is straightforward to implement—just iterate through all 128 features, multiply corresponding elements from the source and destination feature vectors, and accumulate the sum. The implementation would use `idx < num_edges` as the guard condition to handle arbitrary edge counts safely. However, I'm reconsidering whether a warp-per-edge approach might be more performant since it enables better memory coalescing across the feature dimension. The trade-off is that each warp only handles a single edge, so I'd need to verify this actually improves throughput given the overhead of coordinating across threads.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n# Define and compile the CUDA kernel\ncuda_src = r'''\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n// Warp-level all-reduce using shuffle\n__device__ __inline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset >>= 1) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Optimized kernel: each warp handles one edge for better memory coalescing\n__global__ void edge_matmul_kernel(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n // Each warp handles one edge\n const int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;\n const int lane_id = threadIdx.x % 32;\n \n // Grid-stride loop over edges\n for (int edge = warp_id; edge < num_edges; edge += gridDim.x * blockDim.x / 32) {\n // Get the node indices for this edge\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n \n // Compute partial dot product\n // Each thread handles feat_dim/32 features (assuming feat_dim divisible by 32)\n float sum = 0.0f;\n \n // Strided access for coalescing: thread k accesses feature k, k+32, k+64, ...\n // This ensures consecutive threads access consecutive memory\n for (int k = lane_id; k < feat_dim; k += 32) {\n float s_val = src_feat[src * feat_dim + k];\n float d_val = dst_feat[dst * feat_dim + k];\n sum += s_val * d_val;\n }\n \n // Reduce across the warp\n sum = warp_reduce_sum(sum);\n \n // Lane 0 writes the result\n if (lane_id == 0) {\n output[edge] = sum;\n }\n }\n}\n\n// Alternative kernel for non-power-of-2 feature dimensions\n__global__ void edge_matmul_general_kernel(\n const int* __restrict__ src_idx,\n const int* __restrict__ dst_idx,\n const float* __restrict__ src_feat,\n const float* __restrict__ dst_feat,\n float* __restrict__ output,\n int num_edges,\n int feat_dim\n) {\n const int tid = blockIdx.x * blockDim.x + threadIdx.x;\n const int num_threads = gridDim.x * blockDim.x;\n \n // Each thread handles one edge\n for (int edge = tid; edge < num_edges; edge += num_threads) {\n int src = src_idx[edge];\n int dst = dst_idx[edge];\n \n float sum = 0.0f;\n for (int k = 0; k < feat_dim; k++) {\n float s_val = src_feat[src * feat_dim + k];\n float d_val = dst_feat[dst * feat_dim + k];\n sum += s_val * d_val;\n }\n \n output[edge] = sum;\n }\n}\n\n// Host wrapper - choose optimal kernel based on problem size\ntorch::Tensor edge_matmul_forward(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n) {\n const int num_edges = src_idx.size(0);\n const int feat_dim = src_feat.size(1);\n \n auto output = torch::empty({num_edges}, src_feat.options());\n \n const int threads_per_block = 128; // 4 warps per block\n \n if (feat_dim % 32 == 0 && feat_dim >= 32) {\n // Use warp-optimized kernel\n // Each warp handles one edge\n const int warps_per_block = threads_per_block / 32;\n const int num_blocks = (num_edges + warps_per_block - 1) / warps_per_block;\n const int max_blocks = 65535;\n const int actual_blocks = min(num_blocks, max_blocks);\n \n edge_matmul_kernel<<<actual_blocks, threads_per_block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n } else {\n // Use general kernel with one thread per edge\n const int num_blocks = (num_edges + threads_per_block - 1) / threads_per_block;\n const int max_blocks = 65535;\n const int actual_blocks = min(num_blocks, max_blocks);\n \n edge_matmul_general_kernel<<<actual_blocks, threads_per_block>>>(\n src_idx.data_ptr<int>(),\n dst_idx.data_ptr<int>(),\n src_feat.data_ptr<float>(),\n dst_feat.data_ptr<float>(),\n output.data_ptr<float>(),\n num_edges,\n feat_dim\n );\n }\n \n return output;\n}\n'''\n\ncpp_src = r'''\ntorch::Tensor edge_matmul_forward(\n torch::Tensor src_idx,\n torch::Tensor dst_idx,\n torch::Tensor src_feat,\n torch::Tensor dst_feat\n);\n'''\n\n# Load the CUDA extension\nedge_matmul_ext = load_inline(\n name='edge_matmul_ext',\n cpp_sources=cpp_src,\n cuda_sources=cuda_src,\n functions=['edge_matmul_forward'],\n extra_cuda_cflags=['-O3', '-arch=sm_80', '--use_fast_math'],\n verbose=False\n)\n\n\nclass ModelNew(nn.Module):\n \"\"\"\n Edge-only sampled dense-dense matmul with custom CUDA kernel.\n \n For each graph edge (src, dst), compute a dot product between source and\n destination node embeddings using optimized warp-level parallelism.\n \"\"\"\n \n def __init__(self):\n super().__init__()\n \n def forward(\n self,\n src_idx: torch.Tensor,\n dst_idx: torch.Tensor,\n src_feat: torch.Tensor,\n dst_feat: torch.Tensor,\n ) -> torch.Tensor:\n return edge_matmul_ext.edge_matmul_forward(\n src_idx, dst_idx, src_feat, dst_feat\n )" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: correctness check did not pass (? trials). runtime_error: expected scalar type Int but found Float
final submitted kernel (181 lines, 5,573 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

# Define and compile the CUDA kernel
cuda_src = r'''
#include <cuda.h>
#include <cuda_runtime.h>

// Warp-level all-reduce using shuffle
__device__ __inline__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}

// Optimized kernel: each warp handles one edge for better memory coalescing
__global__ void edge_matmul_kernel(
    const int* __restrict__ src_idx,
    const int* __restrict__ dst_idx,
    const float* __restrict__ src_feat,
    const float* __restrict__ dst_feat,
    float* __restrict__ output,
    int num_edges,
    int feat_dim
) {
    // Each warp handles one edge
    const int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
    const int lane_id = threadIdx.x % 32;
    
    // Grid-stride loop over edges
    for (int edge = warp_id; edge < num_edges; edge += gridDim.x * blockDim.x / 32) {
        // Get the node indices for this edge
        int src = src_idx[edge];
        int dst = dst_idx[edge];
        
        // Compute partial dot product
        // Each thread handles feat_dim/32 features (assuming feat_dim divisible by 32)
        float sum = 0.0f;
        
        // Strided access for coalescing: thread k accesses feature k, k+32, k+64, ...
        // This ensures consecutive threads access consecutive memory
        for (int k = lane_id; k < feat_dim; k += 32) {
            float s_val = src_feat[src * feat_dim + k];
            float d_val = dst_feat[dst * feat_dim + k];
            sum += s_val * d_val;
        }
        
        // Reduce across the warp
        sum = warp_reduce_sum(sum);
        
        // Lane 0 writes the result
        if (lane_id == 0) {
            output[edge] = sum;
        }
    }
}

// Alternative kernel for non-power-of-2 feature dimensions
__global__ void edge_matmul_general_kernel(
    const int* __restrict__ src_idx,
    const int* __restrict__ dst_idx,
    const float* __restrict__ src_feat,
    const float* __restrict__ dst_feat,
    float* __restrict__ output,
    int num_edges,
    int feat_dim
) {
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int num_threads = gridDim.x * blockDim.x;
    
    // Each thread handles one edge
    for (int edge = tid; edge < num_edges; edge += num_threads) {
        int src = src_idx[edge];
        int dst = dst_idx[edge];
        
        float sum = 0.0f;
        for (int k = 0; k < feat_dim; k++) {
            float s_val = src_feat[src * feat_dim + k];
            float d_val = dst_feat[dst * feat_dim + k];
            sum += s_val * d_val;
        }
        
        output[edge] = sum;
    }
}

// Host wrapper - choose optimal kernel based on problem size
torch::Tensor edge_matmul_forward(
    torch::Tensor src_idx,
    torch::Tensor dst_idx,
    torch::Tensor src_feat,
    torch::Tensor dst_feat
) {
    const int num_edges = src_idx.size(0);
    const int feat_dim = src_feat.size(1);
    
    auto output = torch::empty({num_edges}, src_feat.options());
    
    const int threads_per_block = 128;  // 4 warps per block
    
    if (feat_dim % 32 == 0 && feat_dim >= 32) {
        // Use warp-optimized kernel
        // Each warp handles one edge
        const int warps_per_block = threads_per_block / 32;
        const int num_blocks = (num_edges + warps_per_block - 1) / warps_per_block;
        const int max_blocks = 65535;
        const int actual_blocks = min(num_blocks, max_blocks);
        
        edge_matmul_kernel<<<actual_blocks, threads_per_block>>>(
            src_idx.data_ptr<int>(),
            dst_idx.data_ptr<int>(),
            src_feat.data_ptr<float>(),
            dst_feat.data_ptr<float>(),
            output.data_ptr<float>(),
            num_edges,
            feat_dim
        );
    } else {
        // Use general kernel with one thread per edge
        const int num_blocks = (num_edges + threads_per_block - 1) / threads_per_block;
        const int max_blocks = 65535;
        const int actual_blocks = min(num_blocks, max_blocks);
        
        edge_matmul_general_kernel<<<actual_blocks, threads_per_block>>>(
            src_idx.data_ptr<int>(),
            dst_idx.data_ptr<int>(),
            src_feat.data_ptr<float>(),
            dst_feat.data_ptr<float>(),
            output.data_ptr<float>(),
            num_edges,
            feat_dim
        );
    }
    
    return output;
}
'''

cpp_src = r'''
torch::Tensor edge_matmul_forward(
    torch::Tensor src_idx,
    torch::Tensor dst_idx,
    torch::Tensor src_feat,
    torch::Tensor dst_feat
);
'''

# Load the CUDA extension
edge_matmul_ext = load_inline(
    name='edge_matmul_ext',
    cpp_sources=cpp_src,
    cuda_sources=cuda_src,
    functions=['edge_matmul_forward'],
    extra_cuda_cflags=['-O3', '-arch=sm_80', '--use_fast_math'],
    verbose=False
)


class ModelNew(nn.Module):
    """
    Edge-only sampled dense-dense matmul with custom CUDA kernel.
    
    For each graph edge (src, dst), compute a dot product between source and
    destination node embeddings using optimized warp-level parallelism.
    """
    
    def __init__(self):
        super().__init__()
    
    def forward(
        self,
        src_idx: torch.Tensor,
        dst_idx: torch.Tensor,
        src_feat: torch.Tensor,
        dst_feat: torch.Tensor,
    ) -> torch.Tensor:
        return edge_matmul_ext.edge_matmul_forward(
            src_idx, dst_idx, src_feat, dst_feat
        )