Kernel Details - modular_device_functions_edit

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,
    multiplier: torch.Tensor,
) -> torch.Tensor:
    """
    Applies convolution, scalar multiplication, LeakyReLU and GELU.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, height, width)
        conv_weight (torch.Tensor): Convolution weights of shape (out_channels, in_channels, kernel_size, kernel_size)
        conv_bias (torch.Tensor): Convolution bias of shape (out_channels)
        multiplier (torch.Tensor): Learnable scalar of shape (out_channels, 1, 1)

    Returns:
        torch.Tensor: Output tensor after applying convolution, multiplication, LeakyReLU and GELU
    """
    x = F.conv2d(x, conv_weight, bias=conv_bias)
    x = x * multiplier
    x = F.leaky_relu(x)
    x = F.gelu(x)
    return x


class Model(nn.Module):
    """
    Model that performs a convolution, multiplies by a learnable scalar, applies LeakyReLU, and then GELU.
    """

    def __init__(self, in_channels, out_channels, kernel_size, multiplier_shape):
        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)
        self.multiplier = nn.Parameter(torch.randn(multiplier_shape) * 0.02)

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


batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3
multiplier_shape = (out_channels, 1, 1)


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


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

import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Model that performs a convolution, multiplies by a learnable scalar, applies LeakyReLU, and then GELU.
    """
    def __init__(self, in_channels, out_channels, kernel_size, multiplier_shape):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.multiplier = nn.Parameter(torch.randn(multiplier_shape) * 0.02) 
        self.leaky_relu = nn.LeakyReLU()

    def forward(self, x):
        x = self.conv(x)
        x = x * self.multiplier
        x = self.leaky_relu(x)
        x = torch.nn.functional.gelu(x)
        return x

batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3
multiplier_shape = (out_channels, 1, 1)

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

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

Download Evaluation Download PyTorch Download CUDA Download Profiles

Kernel Information

Operation Name	54_Conv2d_Multiply_LeakyReLU_GELU
Level ID	2
Task ID	54
Kernel Name	modular_device_functions_edit_1
CUDA Speedup (Native)	1.249x
CUDA Speedup (Compile)	1.404x
CUDA Runtime	0.040 ms
PyTorch Runtime (Native)	0.050 ms
PyTorch Runtime (Compile)	0.056 ms
Correct	True
Max Diff (vs. Reference)	0.000000
Model	azure-gpt-4o-2024-08-06
Temperature	0.00

View Experiment Progress Details

Related Kernels (Level 2, Task 54 • 54_Conv2d_Multiply_LeakyReLU_GELU)

Rank	Kernel Name	Runtime (ms)	Speedup Native	Speedup Compile
🥇	54_Conv2d_Multiply_LeakyReLU_GELU	0.04	1.28	1.44
🥇	balanced_workload_distribution_base	0.04	1.28	1.44
🥇	warp_divergence_optimized_base	0.04	1.28	1.44
🥇	optimized_block_size_128_base	0.04	1.28	1.44
🥇	optimized_convolution_with_tunable_blocksize_base	0.04	1.28	1.44
🥇	direct_3d_indexing_opt_base	0.04	1.28	1.44
🥇	direct_3d_indexing_base	0.04	1.28	1.44
🥇	unroll_loops_54conv_edit_1	0.04	1.28	1.44
🥇	dynamic_block_size_54conv_base	0.04	1.28	1.44
🥇	threadblock_3d_mapping_base	0.04	1.28	1.44
🥇	balanced_thread_distribution_base	0.04	1.28	1.44
🥇	branchless_no_divergence_54conv_base	0.04	1.28	1.44
🥇	modular_device_functions_base	0.04	1.28	1.44
🥇	tile_based_2d_indexing_base	0.04	1.28	1.44
15	combined_conv_act_base	0.04	1.25	1.40
15	optimized_stride_loop_base	0.04	1.25	1.40
15	unroll_loops_54conv_base	0.04	1.25	1.40
15	54_Conv2d_Multiply_LeakyReLU_GELU_warp_divergence_reduction_base	0.04	1.25	1.40
15	dynamic_block_size_54conv_edit_1	0.04	1.25	1.40
15	modular_device_functions_edit_1	0.04	1.25	1.40

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

// Device function: GELU approximation
__device__ __forceinline__ float gelu(float x) {
    const float k0 = 0.7978845608028654f; // sqrt(2/pi)
    return 0.5f * x * (1.0f + tanhf(k0 * (x + 0.044715f * x * x * x)));
}

// Device function: LeakyReLU activation
__device__ __forceinline__ float leaky_relu(float x) {
    return fmaxf(x, 0.01f * x);
}

// Device function: Convolution operation
__device__ float convolution(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    int n, int oc, int oh, int ow,
    int in_channels, int input_h, int input_w,
    int kernel_size
) {
    float sum = 0.0f;
    for (int ic = 0; ic < in_channels; ic++) {
        for (int i = 0; i < kernel_size; i++) {
            for (int j = 0; j < kernel_size; j++) {
                int in_h = oh + i; // stride = 1, no padding. if (in_h >= input_h) continue;
                int in_w = ow + j;
                int input_index = ((n * in_channels + ic) * input_h + in_h) * input_w + in_w;
                int weight_index = ((oc * in_channels + ic) * kernel_size + i) * kernel_size + j;
                sum += input[input_index] * weight[weight_index];
            }
        }
    }
    return sum;
}

// CUDA kernel that performs convolution, scalar multiplication, LeakyReLU and GELU.
__global__ void conv_forward_kernel(
    const float* __restrict__ input,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    const float* __restrict__ multiplier,
    float* __restrict__ output,
    int batch_size,
    int in_channels,
    int input_h,
    int input_w,
    int out_channels,
    int kernel_size,
    int output_h,
    int output_w
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = batch_size * out_channels * output_h * output_w;
    int stride = blockDim.x * gridDim.x;
    
    // Grid-stride loop to cover all output elements.
    for (int index = idx; index < total; index += stride) {
        // Calculate indices for output
        int ow = index % output_w;
        int tmp = index / output_w;
        int oh = tmp % output_h;
        tmp = tmp / output_h;
        int oc = tmp % out_channels;
        int n = tmp / out_channels;

        // Start with the bias for output channel oc.
        float sum = bias[oc];
        
        // Perform convolution
        sum += convolution(input, weight, n, oc, oh, ow, in_channels, input_h, input_w, kernel_size);
        
        // Multiply with the channel-specific multiplier.
        sum *= multiplier[oc];
        
        // Apply LeakyReLU activation.
        sum = leaky_relu(sum);
        
        // Apply GELU activation.
        output[index] = gelu(sum);
    }
}

// C++ interface (to be called from Python)
torch::Tensor forward_cuda(
    torch::Tensor input,
    torch::Tensor conv_weight,
    torch::Tensor conv_bias,
    torch::Tensor multiplier
) {
    // Get input dimensions.
    const auto batch_size = input.size(0);
    const auto in_channels = input.size(1);
    const auto input_h = input.size(2);
    const auto input_w = input.size(3);
    
    // Get convolution parameters.
    const auto out_channels = conv_weight.size(0);
    const auto kernel_size = conv_weight.size(2);
    const auto output_h = input_h - kernel_size + 1;
    const auto output_w = input_w - kernel_size + 1;
    
    // Allocate output tensor.
    auto output = torch::empty({batch_size, out_channels, output_h, output_w}, input.options());
    
    // Launch CUDA kernel.
    const int total_elements = batch_size * out_channels * output_h * output_w;
    const int threads = 256;
    const int blocks = (total_elements + threads - 1) / threads;
    
    conv_forward_kernel<<<blocks, threads>>>(
        input.data_ptr<float>(),
        conv_weight.data_ptr<float>(),
        conv_bias.data_ptr<float>(),
        multiplier.data_ptr<float>(),
        output.data_ptr<float>(),
        batch_size,
        in_channels,
        input_h,
        input_w,
        out_channels,
        kernel_size,
        output_h,
        output_w
    );
    
    // Check for kernel errors.
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("CUDA kernel failed: %s\n", cudaGetErrorString(err));
    }
    
    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward_cuda, "Convolution, scalar multiplication, LeakyReLU and GELU (CUDA)");
}

Performance Metrics

Metric	Value	Unit	Variance	Samples
Executed Ipc Active	3.186	inst/cycle	0.000	5
Executed Ipc Elapsed	2.960	inst/cycle	0.000	5
Issue Slots Busy	79.894	%	0.024	5
Issued Ipc Active	3.196	inst/cycle	0.000	5
SM Busy	79.894	%	0.024	5
Memory Throughput	35280253588.192	byte/second	12779322814515072.000	5
Mem Busy	52.950	%	0.015	5
Max Bandwidth	36.182	%	0.007	5
L1/TEX Hit Rate	87.926	%	0.000	5
L2 Hit Rate	92.754	%	0.027	5
Mem Pipes Busy	35.666	%	0.007	5
Warp Cycles Per Issued Instruction	16.650	cycle	0.000	5
Warp Cycles Per Executed Instruction	16.690	cycle	0.000	5
Avg. Active Threads Per Warp	32.000		0.000	5
Avg. Not Predicated Off Threads Per Warp	29.360		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	83.580	%	0.004	5
Achieved Active Warps Per SM	53.490	warp	0.001	5

Analysis Rules

Rule	Description
INF HighPipeUtilization	ALU is the highest-utilized pipeline (48.1%) 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.
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 (83.5%) 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.

Rule

Description

INF HighPipeUtilization

ALU is the highest-utilized pipeline (48.1%) 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.

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 (83.5%) 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	205748.76	μs
Device Time	85.73	μs
Self CPU Time	58.76	μ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	205690.00	μs
Device Time	85.73	μs
Self CPU Time	126.88	μ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	671379.85	μs
Device Time	15120.76	μs
Self CPU Time	671379.85	μs
Self Device Time	15120.76	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
conv_forward_kernel(float const, float const, float const, float const, float*, int, int, int, int, int, int, int, int)
CPU Time	0.00	μs
Device Time	272557.85	μs
Self CPU Time	0.00	μs
Self Device Time	272557.85	μ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	20451.02	μs
Device Time	30226.19	μs
Self CPU Time	20451.02	μs
Self Device Time	30226.19	μ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	236415.77	μs
Device Time	582795.57	μs
Self CPU Time	13359.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	223058.48	μs
Device Time	582795.57	μs
Self CPU Time	17230.24	μs
Self Device Time	582795.57	μ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	582873.84	μs
Self CPU Time	0.00	μs
Self Device Time	582873.84	μ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

45299 warnings generated when compiling for host.
Suppressed 45324 warnings (45277 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/20250212_optimize_b5_s4_e1_v2/level_2/task_54/b3_s0_modular_device_functions/edit_1/edit_1.cu:23:5 bugprone-easily-swappable-parameters