← Back to Leaderboard

The AI CUDA Engineer 👷

67_Conv2d_GELU_GlobalAvgPoolwarp_optimized_conv2d_gelu_pool_base

Level 2 • Task 67
import torch
import torch.nn as nn
import torch.nn.functional as F


def module_fn(
    x: torch.Tensor,
    conv_weight: torch.Tensor,
    conv_bias: torch.Tensor,
) -> torch.Tensor:
    """
    Applies convolution, GELU activation, and global average pooling.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, height, width)
        conv_weight (torch.Tensor): Convolution weight tensor of shape
            (out_channels, in_channels, kernel_size, kernel_size)
        conv_bias (torch.Tensor): Convolution bias tensor of shape (out_channels)

    Returns:
        torch.Tensor: Output tensor of shape (batch_size, out_channels)
    """
    x = F.conv2d(x, conv_weight, bias=conv_bias)
    x = F.gelu(x)
    x = F.adaptive_avg_pool2d(x, 1)
    x = x.squeeze(-1).squeeze(-1)
    return x


class Model(nn.Module):
    """
    Simple model that performs a convolution, applies GELU, and then performs global average pooling.
    """

    def __init__(self, in_channels, out_channels, kernel_size):
        super(Model, self).__init__()
        conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.conv_weight = nn.Parameter(conv.weight)
        self.conv_bias = nn.Parameter(conv.bias)

    def forward(self, x, fn=module_fn):
        return fn(x, self.conv_weight, self.conv_bias)


batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3


def get_inputs():
    return [torch.randn(batch_size, in_channels, height, width)]


def get_init_inputs():
    return [in_channels, out_channels, kernel_size]
import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs a convolution, applies GELU, and then performs global average pooling.
    """
    def __init__(self, in_channels, out_channels, kernel_size):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)

    def forward(self, x):
        """
        Args:
            x: Input tensor of shape (batch_size, in_channels, height, width)
        Returns:
            Output tensor of shape (batch_size, out_channels)
        """
        x = self.conv(x)
        x = torch.nn.functional.gelu(x)
        x = torch.nn.functional.adaptive_avg_pool2d(x, 1)
        x = x.squeeze(-1).squeeze(-1)
        return x

batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3

def get_inputs():
    return [torch.randn(batch_size, in_channels, height, width)]

def get_init_inputs():
    return [in_channels, out_channels, kernel_size]

Kernel Information

Related Kernels (Level 2, Task 67 • 67_Conv2d_GELU_GlobalAvgPool)

#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>

#define KERNEL_SIZE 3

// GELU activation function
__device__ inline float gelu_activate(float x) {
    return 0.5f * x * (1.f + erff(x / 1.41421356f));
}

// Convolution computation with unrolling for a fixed kernel size
template<int ksize>
__device__ inline float compute_conv_unrolled(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const int in_channels,
    const int in_h,
    const int in_w,
    const int n,
    const int c_out,
    const int out_row,
    const int out_col
) {
    float sum = 0.0f;
    #pragma unroll
    for (int c_in = 0; c_in < in_channels; c_in++) {
        #pragma unroll
        for (int kh = 0; kh < ksize; kh++) {
            #pragma unroll
            for (int kw = 0; kw < ksize; kw++) {
                int in_row = out_row + kh;
                int in_col = out_col + kw;
                int input_idx = ((n * in_channels + c_in) * in_h + in_row) * in_w + in_col;
                int weight_idx = ((c_out * in_channels + c_in) * ksize + kh) * ksize + kw;
                sum += input[input_idx] * weight[weight_idx];
            }
        }
    }
    return sum;
}

// Warp-level reduction using shuffle
__device__ inline float warp_reduce_sum(float val) {
    // Assuming warp size of 32
    for (int offset = 16; offset > 0; offset /= 2)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;
}

// Fused kernel: performs convolution, applies bias, GELU activation, and accumulates for global average pooling
// Each block processes one (n, c_out) pair over the conv output (spatial dimensions out_h x out_w).
extern "C" __global__ void conv2d_gelu_pool_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ output,
    const int N,
    const int in_channels,
    const int in_h,
    const int in_w,
    const int out_channels,
    const int out_h,
    const int out_w
) {
    // Grid: blockIdx.y -> batch index (n), blockIdx.x -> output channel (c_out)
    const int n = blockIdx.y;
    const int c_out = blockIdx.x;
    
    const int total_pixels = out_h * out_w;
    float local_sum = 0.0f;

    // Each thread processes a subset of the spatial positions in a grid-stride loop
    for (int idx = threadIdx.x; idx < total_pixels; idx += blockDim.x) {
        int row = idx / out_w;
        int col = idx % out_w;
        float conv_val = compute_conv_unrolled<KERNEL_SIZE>(
            input, weight, in_channels, in_h, in_w,
            n, c_out, row, col
        );
        conv_val += bias[c_out];
        conv_val = gelu_activate(conv_val);
        local_sum += conv_val;
    }

    // Perform warp-level reduction
    local_sum = warp_reduce_sum(local_sum);

    // Use shared memory to reduce across warps in the block
    if (threadIdx.x % 32 == 0) {
        atomicAdd(&output[n * out_channels + c_out], local_sum / float(total_pixels));
    }
}

// Host function that launches the fused kernel
torch::Tensor forward(
    torch::Tensor input,
    torch::Tensor conv_weight,
    torch::Tensor conv_bias
) {
    TORCH_CHECK(input.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(conv_weight.is_cuda(), "conv_weight must be a CUDA tensor");
    TORCH_CHECK(conv_bias.is_cuda(), "conv_bias must be a CUDA tensor");

    const int N = input.size(0);
    const int in_channels = input.size(1);
    const int in_h = input.size(2);
    const int in_w = input.size(3);
    const int out_channels = conv_weight.size(0);
    // For a 3x3 filter, output dimensions are (in_h - 2, in_w - 2)
    const int out_h = in_h - 2;
    const int out_w = in_w - 2;

    auto options = torch::TensorOptions().dtype(input.dtype()).device(input.device());
    // Final output has shape (N, out_channels) after global average pooling
    auto final_output = torch::zeros({N, out_channels}, options);

    // Launch the fused kernel with a 2D grid: (c_out, N)
    dim3 grid(out_channels, N);
    const int threads = 256;

    conv2d_gelu_pool_kernel<<<grid, threads>>>(
        input.data_ptr<float>(),
        conv_weight.data_ptr<float>(),
        conv_bias.data_ptr<float>(),
        final_output.data_ptr<float>(),
        N, in_channels, in_h, in_w,
        out_channels, out_h, out_w
    );

    return final_output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward, "Warp-optimized Conv2d + GELU + GlobalAvgPool");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 1.866 inst/cycle 0.000 5
Executed Ipc Elapsed 1.594 inst/cycle 0.001 5
Issue Slots Busy 46.780 % 0.091 5
Issued Ipc Active 1.872 inst/cycle 0.000 5
SM Busy 46.780 % 0.091 5
Memory Throughput 44554574796.030 byte/second 821613541537177856.000 5
Mem Busy 66.226 % 1.825 5
Max Bandwidth 51.134 % 1.079 5
L1/TEX Hit Rate 92.078 % 0.002 5
L2 Hit Rate 84.616 % 2.517 5
Mem Pipes Busy 45.162 % 0.852 5
Warp Cycles Per Issued Instruction 27.518 cycle 0.011 5
Warp Cycles Per Executed Instruction 27.620 cycle 0.010 5
Avg. Active Threads Per Warp 30.670 0.000 5
Avg. Not Predicated Off Threads Per Warp 29.190 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 8.000 block 0.000 5
Block Limit Shared Mem 32.000 block 0.000 5
Block Limit Warps 8.000 block 0.000 5
Theoretical Active Warps per SM 64.000 warp 0.000 5
Theoretical Occupancy 100.000 % 0.000 5
Achieved Occupancy 80.258 % 0.211 5
Achieved Active Warps Per SM 51.362 warp 0.086 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (23.3%) based on active cycles, taking into account the rates of its different instructions. It executes integer and logic operations. It is well-utilized, but should not be a bottleneck.
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 is not impacted by any block limit. The difference between calculated theoretical (100.0%) and measured achieved occupancy (80.8%) can be the result of warp scheduling overheads or workload imbalances during the kernel execution. Load imbalances can occur between warps within a block as well as across blocks of the same kernel. 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 380010.12 μs
Device Time 80.51 μs
Self CPU Time 79.78 μ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::zeros
CPU Time 5381517.37 μs
Device Time 188571.28 μs
Self CPU Time 118642.30 μ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::zero_
CPU Time 6653533.58 μs
Device Time 6337113.50 μs
Self CPU Time 243464.45 μ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 6410071.03 μs
Device Time 6337113.50 μs
Self CPU Time 332463.14 μs
Self Device Time 6337110.81 μ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 6367248.64 μs
Device Time 2518.61 μs
Self CPU Time 6367248.64 μs
Self Device Time 2518.61 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
conv2d_gelu_pool_kernel
CPU Time 0.00 μs
Device Time 2138231.06 μs
Self CPU Time 0.00 μs
Self Device Time 2138231.06 μ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 173697.76 μs
Device Time 1019327.09 μs
Self CPU Time 173697.76 μs
Self Device Time 1019327.09 μ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 6148542.21 μs
Self CPU Time 0.00 μs
Self Device Time 6148542.21 μ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
45292 warnings generated when compiling for host.
Suppressed 45323 warnings (45276 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/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:21:5 bugprone-easily-swappable-parameters
21 | const int n,
| ^~~~~~~~~~~~
22 | const int c_out,
| ~~~~~~~~~~~~~~~~
23 | const int out_row,
| ~~~~~~~~~~~~~~~~~~
24 | const int out_col
| ~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:21:15: note: the first parameter in the range is 'n'
21 | const int n,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:24:15: note: the last parameter in the range is 'out_col'
24 | const int out_col
| ^~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:56:5: warning: 2 adjacent parameters of 'conv2d_gelu_pool_kernel' of similar type ('const float *__restrict') are easily swapped by mistake [bugprone-easily-swappable-parameters]
56 | const float* __restrict__ weight,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
57 | const float* __restrict__ bias,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:56:31: note: the first parameter in the range is 'weight'
56 | const float* __restrict__ weight,
| ^~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:57:31: note: the last parameter in the range is 'bias'
57 | const float* __restrict__ bias,
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:59:5: warning: 2 adjacent parameters of 'conv2d_gelu_pool_kernel' of similar type ('const int') are easily swapped by mistake [bugprone-easily-swappable-parameters]
59 | const int N,
| ^~~~~~~~~~~~
60 | const int in_channels,
| ~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:59:15: note: the first parameter in the range is 'N'
59 | const int N,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:60:15: note: the last parameter in the range is 'in_channels'
60 | const int in_channels,
| ^~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:62:5: warning: 3 adjacent parameters of 'conv2d_gelu_pool_kernel' of similar type ('const int') are easily swapped by mistake [bugprone-easily-swappable-parameters]
62 | const int in_w,
| ^~~~~~~~~~~~~~~
63 | const int out_channels,
| ~~~~~~~~~~~~~~~~~~~~~~~
64 | const int out_h,
| ~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:62:15: note: the first parameter in the range is 'in_w'
62 | const int in_w,
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:64:15: note: the last parameter in the range is 'out_h'
64 | const int out_h,
| ^~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:68:19: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
68 | const int n = blockIdx.y;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:69:23: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
69 | const int c_out = blockIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:75:20: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
75 | for (int idx = threadIdx.x; idx < total_pixels; idx += blockDim.x) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:75:60: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
75 | for (int idx = threadIdx.x; idx < total_pixels; idx += blockDim.x) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:98:19: warning: the parameter 'input' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
98 | torch::Tensor input,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:99:19: warning: the parameter 'conv_weight' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
99 | torch::Tensor conv_weight,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:100:19: warning: the parameter 'conv_bias' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
100 | torch::Tensor conv_bias
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:106:19: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
106 | const int N = input.size(0);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:107:29: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
107 | const int in_channels = input.size(1);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:108:22: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
108 | const int in_h = input.size(2);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:109:22: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
109 | const int in_w = input.size(3);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_67/b5_s1_warp_optimized_conv2d_gelu_pool/base/base.cu:110:30: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
110 | const int out_channels = conv_weight.size(0);
| ^