← Back to Leaderboard

The AI CUDA Engineer 👷

47_Sum_reduction_over_a_dimension • grid_strided_warp_reduction_base

Level 1 • Task 47
import torch
import torch.nn as nn
import torch.nn.functional as F


def module_fn(x: torch.Tensor, dim: int) -> torch.Tensor:
    """
    Applies sum reduction over the specified dimension.

    Args:
        x (torch.Tensor): Input tensor of shape (..., dim, ...).
        dim (int): Dimension to reduce over.

    Returns:
        torch.Tensor: Output tensor after sum reduction, shape (..., 1, ...).
    """
    return torch.sum(x, dim=dim, keepdim=True)


class Model(nn.Module):
    """
    Simple model that performs sum reduction over a specified dimension.
    """

    def __init__(self, dim: int):
        """
        Initializes the model with the dimension to reduce over.

        Args:
            dim (int): Dimension to reduce over.
        """
        super(Model, self).__init__()
        self.dim = dim

    def forward(self, x: torch.Tensor, fn=module_fn) -> torch.Tensor:
        """
        Applies sum reduction over the specified dimension.

        Args:
            x (torch.Tensor): Input tensor of shape (..., dim, ...).

        Returns:
            torch.Tensor: Output tensor after sum reduction, shape (..., 1, ...).
        """
        return fn(x, self.dim)


batch_size = 16
dim1 = 256
dim2 = 256
reduce_dim = 1


def get_inputs():
    x = torch.randn(batch_size, dim1, dim2)
    return [x]


def get_init_inputs():
    return [reduce_dim]

import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs sum reduction over a specified dimension.
    """
    def __init__(self, dim: int):
        """
        Initializes the model with the dimension to reduce over.

        Args:
            dim (int): Dimension to reduce over.
        """
        super(Model, self).__init__()
        self.dim = dim

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        """
        Applies sum reduction over the specified dimension.

        Args:
            x (torch.Tensor): Input tensor of shape (..., dim, ...).

        Returns:
            torch.Tensor: Output tensor after sum reduction, shape (..., 1, ...).
        """
        return torch.sum(x, dim=self.dim, keepdim=True)

batch_size = 16
dim1 = 256
dim2 = 256
reduce_dim = 1

def get_inputs():
    x = torch.randn(batch_size, dim1, dim2)
    return [x]

def get_init_inputs():
    return [reduce_dim]
Download Evaluation Download PyTorch Download CUDA Download Profiles

Kernel Information

Operation Name 47_Sum_reduction_over_a_dimension
Level ID 1
Task ID 47
Kernel Name grid_strided_warp_reduction_base
CUDA Speedup (Native) 1.155x
CUDA Speedup (Compile) 2.333x
CUDA Runtime 0.014 ms
PyTorch Runtime (Native) 0.016 ms
PyTorch Runtime (Compile) 0.033 ms
Correct True
Max Diff (vs. Reference) 0.000000
Model bedrock/anthropic.claude-3-5-sonnet-20241022-v2:0
Temperature 0.50
View Experiment Progress Details

Related Kernels (Level 1, Task 47 • 47_Sum_reduction_over_a_dimension)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 unroll_warp_level_sum_reduction_base 0.01 1.47 2.97
🥇 adaptive_blocksize_warp_sum_reduction_base 0.01 1.47 2.97
🥇 modular_warp_sum_reduction_base 0.01 1.47 2.97
🥇 fully_unrolled_warp_sum_reduction_base 0.01 1.47 2.97
5 tuned_warp_block_size_sum_reduction_base 0.01 1.35 2.72
5 warp_level_primitives_sum_reduction_base 0.01 1.35 2.72
5 uniform_control_flow_warp_sum_reduction_base_base 0.01 1.35 2.72
5 manual_unroll_warp_sum_reduction_base 0.01 1.35 2.72
9 hybrid_warp_shared_reduce_base 0.01 1.16 2.33
9 block_shared_sum_reduction_min_sync_base 0.01 1.16 2.33
9 warp_sum_reduce_edit_1 0.01 1.16 2.33
9 warp_sum_reduce_base 0.01 1.16 2.33
9 warp_sum_reduce_unroll_base 0.01 1.16 2.33
9 warp_sum_reduce_optimized_sync_base 0.01 1.16 2.33
9 warp_shuffle_sum_reduction_optimized_edit_1 0.01 1.16 2.33
9 warp_level_sum_reduction_edit_1 0.01 1.16 2.33
9 cg_warp_uniform_sum_reduction_base 0.01 1.16 2.33
9 cg_warp_uniform_sum_reduction_edit_1 0.01 1.16 2.33
9 grid_strided_warp_reduction_base 0.01 1.16 2.33
9 warp_shuffle_uniform_reduction_edit_1 0.01 1.16 2.33 
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>

template <typename scalar_t>
__global__ void grid_strided_reduce_kernel(
    const scalar_t* __restrict__ input,
    scalar_t* __restrict__ output,
    int64_t reduce_size,
    int64_t outer_size,
    int64_t inner_size) {

    const int tid = threadIdx.x;
    const int bid_x = blockIdx.x;  // for inner dimension
    const int bid_y = blockIdx.y;  // for outer dimension
    const unsigned int FULL_MASK = 0xffffffff;
    
    for (int outer_idx = bid_y; outer_idx < outer_size; outer_idx += gridDim.y) {
        for (int inner_idx = bid_x; inner_idx < inner_size; inner_idx += gridDim.x) {
            const int out_idx = outer_idx * inner_size + inner_idx;
            const int64_t base = outer_idx * reduce_size * inner_size + inner_idx;
            
            scalar_t sum = 0;
            #pragma unroll 4
            for (int i = tid; i < reduce_size; i += blockDim.x) {
                sum += input[base + i * inner_size];
            }
            
            #pragma unroll
            for (int offset = 16; offset > 0; offset >>= 1) {
                sum += __shfl_down_sync(FULL_MASK, sum, offset);
            }
            
            if (tid == 0) {
                output[out_idx] = sum;
            }
        }
    }
}

torch::Tensor grid_strided_reduce_cuda(torch::Tensor input, int64_t dim) {
    if (dim < 0) dim += input.dim();
    
    auto sizes = input.sizes().vec();
    int64_t reduce_size = sizes[dim];
    int64_t outer_size = 1;
    for (int i = 0; i < dim; i++) outer_size *= sizes[i];
    int64_t inner_size = 1;
    for (int i = dim + 1; i < sizes.size(); i++) inner_size *= sizes[i];
    
    sizes[dim] = 1;
    auto output = torch::empty(sizes, input.options());
    
    const int threads = 32;
    dim3 blocks;
    blocks.x = std::min(inner_size, static_cast<int64_t>(256));
    blocks.y = std::min(outer_size, static_cast<int64_t>(256));
    
    AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "grid_strided_reduce_cuda", ([&] {
        grid_strided_reduce_kernel<scalar_t><<<blocks, threads>>>(
            input.data_ptr<scalar_t>(),
            output.data_ptr<scalar_t>(),
            reduce_size,
            outer_size,
            inner_size
        );
    }));
    
    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &grid_strided_reduce_cuda, "Grid-strided reduction (CUDA)");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 0.428 inst/cycle 0.000 5
Executed Ipc Elapsed 0.306 inst/cycle 0.000 5
Issue Slots Busy 10.836 % 0.047 5
Issued Ipc Active 0.432 inst/cycle 0.000 5
SM Busy 10.836 % 0.047 5
Memory Throughput 371757417933.106 byte/second 23165409011895128064.000 5
Mem Busy 56.882 % 0.580 5
Max Bandwidth 32.122 % 0.101 5
L1/TEX Hit Rate 0.000 % 0.000 5
L2 Hit Rate 83.224 % 1.675 5
Mem Pipes Busy 3.850 % 0.003 5
Warp Cycles Per Issued Instruction 63.966 cycle 0.066 5
Warp Cycles Per Executed Instruction 64.816 cycle 0.081 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 28.920 0.000 5
Max Active Clusters 0.000 cluster 0.000 5
Max Cluster Size 8.000 block 0.000 5
Overall GPU Occupancy 0.000 % 0.000 5
Cluster Occupancy 0.000 % 0.000 5
Block Limit SM 32.000 block 0.000 5
Block Limit Registers 84.000 block 0.000 5
Block Limit Shared Mem 32.000 block 0.000 5
Block Limit Warps 64.000 block 0.000 5
Theoretical Active Warps per SM 32.000 warp 0.000 5
Theoretical Occupancy 50.000 % 0.000 5
Achieved Occupancy 44.160 % 0.054 5
Achieved Active Warps Per SM 28.262 warp 0.023 5
Analysis Rules
Rule Description
WRN HighPipeUtilization All compute pipelines are under-utilized. Either this kernel is very small or it doesn't issue enough warps per scheduler. Check the Launch Statistics and Scheduler Statistics sections for further details.
INF CPIStall Check the Warp Stall Sampling (All Cycles) table for the top stall locations in your source based on sampling data. The Kernel Profiling Guide (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-reference) provides more details on each stall reason.
WRN Occupancy This kernel's theoretical occupancy (50.0%) is limited by the number of blocks that can fit on the SM. This kernel's theoretical occupancy (50.0%) is limited by the required amount of shared memory. See the CUDA Best Practices Guide (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy) for more details on optimizing occupancy.
Operation / Metric Value Unit
aten::to
CPU Time 450044.26 μs
Device Time 375.97 μs
Self CPU Time 37.00 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::_to_copy
CPU Time 450007.26 μs
Device Time 375.97 μs
Self CPU Time 95.74 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::empty_strided
CPU Time 449301.16 μs
Device Time 0.00 μs
Self CPU Time 72.00 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaDeviceGetStreamPriorityRange
CPU Time 449036.49 μs
Device Time 0.00 μs
Self CPU Time 449036.49 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaLaunchKernel
CPU Time 443728.31 μs
Device Time 17230.35 μs
Self CPU Time 443728.31 μs
Self Device Time 17230.35 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
void grid_strided_reduce_kernel<float>(float const*, float*, long, long, long)
CPU Time 0.00 μs
Device Time 68721.71 μs
Self CPU Time 0.00 μs
Self Device Time 68721.71 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaEventRecord
CPU Time 14989.07 μs
Device Time 34370.44 μs
Self CPU Time 14989.07 μs
Self Device Time 34370.44 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::zero_
CPU Time 56723.42 μs
Device Time 515269.43 μs
Self CPU Time 10266.17 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::fill_
CPU Time 46459.05 μs
Device Time 515269.43 μs
Self CPU Time 14299.88 μs
Self Device Time 515269.43 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, at::detail::Array<char*, 1> >(int, at::native::FillFunctor<int>, at::detail::Array<char*, 1>)
CPU Time 0.00 μs
Device Time 515269.43 μs
Self CPU Time 0.00 μs
Self Device Time 515269.43 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
Status: Completed 
45287 warnings generated when compiling for host.
Suppressed 45322 warnings (45275 in non-user code, 47 NOLINT).
Use -header-filter=.* to display errors from all non-system headers. Use -system-headers to display errors from system headers as well.
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:9:5 bugprone-easily-swappable-parameters
9 | int64_t reduce_size,
| ^~~~~~~~~~~~~~~~~~~~
10 | int64_t outer_size,
| ~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:9:13: note: the first parameter in the range is 'reduce_size'
9 | int64_t reduce_size,
| ^~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:10:13: note: the last parameter in the range is 'outer_size'
10 | int64_t outer_size,
| ^~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:13:21: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
13 | const int tid = threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:14:23: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
14 | const int bid_x = blockIdx.x; // for inner dimension
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:15:23: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
15 | const int bid_y = blockIdx.y; // for outer dimension
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:18:70: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
18 | for (int outer_idx = bid_y; outer_idx < outer_size; outer_idx += gridDim.y) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:19:74: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
19 | for (int inner_idx = bid_x; inner_idx < inner_size; inner_idx += gridDim.x) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:20:33: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
20 | const int out_idx = outer_idx * inner_size + inner_idx;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:25:53: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
25 | for (int i = tid; i < reduce_size; i += blockDim.x) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:49:18: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
49 | for (int i = dim + 1; i < sizes.size(); i++) inner_size *= sizes[i];
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_1/task_47/b5_s0_grid_strided_warp_reduction/base/base.cu:59:5: warning: inside a lambda, '__func__' expands to the name of the function call operator; consider capturing the name of the enclosing function explicitly [bugprone-lambda-function-name]
59 | AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "grid_strided_reduce_cuda", ([&] {
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:237:34: note: expanded from macro 'AT_DISPATCH_FLOATING_TYPES'
237 | AT_DISPATCH_SWITCH(TYPE, NAME, AT_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:233:3: note: expanded from macro 'AT_DISPATCH_CASE_FLOATING_TYPES'
233 | AT_DISPATCH_CASE(at::ScalarType::Double, __VA_ARGS__) \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:74:3: note: expanded from macro 'AT_DISPATCH_CASE'
74 | AT_PRIVATE_CASE_TYPE_USING_HINT(enum_type, scalar_t, __VA_ARGS__)
| ^
note: (skipping 1 expansions in backtrace; use -fmacro-backtrace-limit=0 to see all)
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:58:7: note: expanded from macro 'AT_PRIVATE_CHECK_SELECTIVE_BUILD'
58 | AT_ERROR( \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/c10/util/Exception.h:711:32: note: expanded from macro 'AT_ERROR'
711 | C10_EXPAND_MSVC_WORKAROUND(TORCH_CHECK(false, ::c10::str(__VA_ARGS__))); \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/c10/util/Exception.h:536:9: note: expanded from macro 'TORCH_CHECK'
536 | __func__, \
| ^