← Back to Leaderboard

The AI CUDA Engineer 👷

48_Conv3d_Scaling_Tanh_Multiply_Sigmoid • unroll_optimization_base

Level 2 • Task 48
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,
    scaling_factor: torch.Tensor,
    bias: torch.Tensor,
) -> torch.Tensor:
    """
    Applies 3D convolution, scaling, tanh, bias multiplication and sigmoid.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, depth, height, width)
        conv_weight (torch.Tensor): 3D convolution weight tensor
        conv_bias (torch.Tensor): 3D convolution bias tensor
        scaling_factor (torch.Tensor): Scaling factor tensor of shape (out_channels, 1, 1, 1)
        bias (torch.Tensor): Bias tensor of shape (out_channels, 1, 1, 1)

    Returns:
        torch.Tensor: Output tensor after applying convolution, scaling, tanh, bias and sigmoid
    """
    x = F.conv3d(x, conv_weight, bias=conv_bias)
    x = x * scaling_factor
    x = torch.tanh(x)
    x = x * bias
    x = torch.sigmoid(x)
    return x


class Model(nn.Module):
    """
    Model that performs a 3D convolution, scales the output, applies tanh, multiplies by a scaling factor, and applies sigmoid.
    """

    def __init__(
        self, in_channels, out_channels, kernel_size, scaling_factor, bias_shape
    ):
        super(Model, self).__init__()
        conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.conv_weight = nn.Parameter(conv.weight)
        self.conv_bias = nn.Parameter(
            conv.bias
            + torch.randn(
                conv.bias.shape, device=conv.bias.device, dtype=conv.bias.dtype
            )
            * 0.02
        )
        self.scaling_factor = nn.Parameter(torch.randn(bias_shape) * 0.02)
        self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)

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


batch_size = 128
in_channels = 3
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
scaling_factor = 2
bias_shape = (out_channels, 1, 1, 1)


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


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

import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Model that performs a 3D convolution, scales the output, applies tanh, multiplies by a scaling factor, and applies sigmoid.
    """
    def __init__(self, in_channels, out_channels, kernel_size, scaling_factor, bias_shape):
        super(Model, self).__init__()
        self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.conv.bias = nn.Parameter(self.conv.bias + torch.randn(self.conv.bias.shape, device=self.conv.bias.device, dtype=self.conv.bias.dtype) * 0.02)
        self.scaling_factor = nn.Parameter(torch.randn(bias_shape) * 0.02)
        self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)

    def forward(self, x):
        x = self.conv(x)
        x = x * self.scaling_factor 
        x = torch.tanh(x)
        x = x * self.bias
        x = torch.sigmoid(x)
        return x

batch_size = 128
in_channels = 3
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
scaling_factor = 2
bias_shape = (out_channels, 1, 1, 1)

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

def get_init_inputs():
    return [in_channels, out_channels, kernel_size, scaling_factor, bias_shape]
Download Evaluation Download PyTorch Download CUDA Download Profiles

Kernel Information

Operation Name 48_Conv3d_Scaling_Tanh_Multiply_Sigmoid
Level ID 2
Task ID 48
Kernel Name unroll_optimization_base
CUDA Speedup (Native) 1.264x
CUDA Speedup (Compile) 0.661x
CUDA Runtime 0.798 ms
PyTorch Runtime (Native) 1.009 ms
PyTorch Runtime (Compile) 0.527 ms
Correct True
Max Diff (vs. Reference) 0.000000
Model azure-gpt-4o-2024-08-06
Temperature 1.00
View Experiment Progress Details

Related Kernels (Level 2, Task 48 • 48_Conv3d_Scaling_Tanh_Multiply_Sigmoid)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 optimized_hybrid_conv3d_base 0.78 1.29 0.68
🥈 streamlined_syncthreads_conv3d_base_edit_1 0.78 1.29 0.67
🥉 aligned_coalesced_conv3d_edit_1 0.78 1.29 0.67
🥉 aligned_memory_access_conv3d_base 0.78 1.29 0.67
🥉 streamlined_syncthreads_conv3d_base_base 0.78 1.29 0.67
6 unrolled_conv3d_opt_edit_1 0.79 1.28 0.67
7 unrolled_conv3d_opt_base 0.79 1.28 0.67
8 block_size_experimentation_edit_1 0.79 1.27 0.67
8 warp_broadcast_tile_edit_1 0.79 1.27 0.67
10 modular_device_functions_edit_1 0.79 1.27 0.67
11 block_size_experimentation_base 0.79 1.27 0.66
12 constant_mem_optimization_base 0.80 1.27 0.66
12 constant_memory_optimization_base 0.80 1.27 0.66
12 constant_mem_opt_base 0.80 1.27 0.66
12 constant_mem_optimization_edit_1 0.80 1.27 0.66
12 strided_loops_conv3d_edit_1 0.80 1.27 0.66
17 constant_memory_optimization_edit_1 0.80 1.27 0.66
18 modular_device_functions_base 0.80 1.26 0.66
18 optimized_shared_mem_edit_1 0.80 1.26 0.66
18 unroll_optimization_base 0.80 1.26 0.66 
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>
#include <math.h>

// Device function for scaling
__device__ float scale_value(float val, float scale) {
    return val * scale;
}

// Device function for tanh
__device__ float apply_tanh(float val) {
    return tanh(val);
}

// Device function for bias multiplication
__device__ float multiply_bias(float val, float bias) {
    return val * bias;
}

// Device function for sigmoid
__device__ float apply_sigmoid(float val) {
    return 1.0f / (1.0f + exp(-val));
}

// CUDA kernel using manual loop unrolling for better performance
__global__ void conv3d_unroll_kernel(
    const float* __restrict__ output,
    const float* __restrict__ scaling_factor,
    const float* __restrict__ bias,
    float* __restrict__ result,
    const int batch_size,
    const int out_channels,
    const int depth,
    const int height, 
    const int width) {
    
    extern __shared__ char smem[];
    float* s_scaling = reinterpret_cast<float*>(smem);
    float* s_bias = s_scaling + out_channels;

    if (threadIdx.x < out_channels) {
        s_scaling[threadIdx.x] = scaling_factor[threadIdx.x];
        s_bias[threadIdx.x] = bias[threadIdx.x];
    }
    __syncthreads();

    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const int total_elements = batch_size * out_channels * depth * height * width;

    if (idx < total_elements) {
        const int c = (idx / (width * height * depth)) % out_channels;

        float val = output[idx];
        val = scale_value(val, s_scaling[c]);
        val = apply_tanh(val);
        val = multiply_bias(val, s_bias[c]);
        val = apply_sigmoid(val);

        result[idx] = val;
    }
}

// Forward function that performs conv3d followed by scaling, tanh, bias multiplication and sigmoid
// It launches the optimized kernel with loop unrolling

torch::Tensor forward(
    torch::Tensor x,
    torch::Tensor conv_weight,
    torch::Tensor conv_bias,
    torch::Tensor scaling_factor,
    torch::Tensor bias) {

    auto conv_out = torch::conv3d(x, conv_weight, conv_bias);
    
    const int batch_size = conv_out.size(0);
    const int out_channels = conv_out.size(1);
    const int depth = conv_out.size(2);
    const int height = conv_out.size(3);
    const int width = conv_out.size(4);

    auto result = torch::empty_like(conv_out);

    const int threads = 256;
    const int total_elements = batch_size * out_channels * depth * height * width;
    const int blocks = (total_elements + threads - 1) / threads;

    conv3d_unroll_kernel<<<blocks, threads, 2 * out_channels * sizeof(float)>>>(
        conv_out.data_ptr<float>(),
        scaling_factor.data_ptr<float>(),
        bias.data_ptr<float>(),
        result.data_ptr<float>(),
        batch_size,
        out_channels, 
        depth,
        height,
        width
    );

    return result;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward, "Conv3d scale tanh bias sigmoid forward with unrolling");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 3.360 inst/cycle 0.000 5
Executed Ipc Elapsed 3.294 inst/cycle 0.000 5
Issue Slots Busy 84.022 % 0.004 5
Issued Ipc Active 3.362 inst/cycle 0.000 5
SM Busy 84.022 % 0.004 5
Memory Throughput 1468562578048.788 byte/second 2065521673129878272.000 5
Mem Busy 28.270 % 0.002 5
Max Bandwidth 43.820 % 0.002 5
L1/TEX Hit Rate 5.870 % 0.000 5
L2 Hit Rate 49.968 % 0.014 5
Mem Pipes Busy 59.424 % 0.007 5
Warp Cycles Per Issued Instruction 15.502 cycle 0.000 5
Warp Cycles Per Executed Instruction 15.502 cycle 0.000 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 25.620 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 16.000 block 0.000 5
Block Limit Shared Mem 28.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 82.266 % 0.002 5
Achieved Active Warps Per SM 52.650 warp 0.001 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (52.9%) 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 (82.3%) 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 727794.15 μs
Device Time 8451.25 μs
Self CPU Time 65.32 μ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 738603.92 μs
Device Time 0.00 μs
Self CPU Time 17931.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::conv3d
CPU Time 627960.09 μs
Device Time 4285903.07 μs
Self CPU Time 11583.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
aten::convolution
CPU Time 616376.60 μs
Device Time 4285903.07 μs
Self CPU Time 15264.98 μ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::_convolution
CPU Time 601111.62 μs
Device Time 4285903.07 μs
Self CPU Time 30407.90 μ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::cudnn_convolution
CPU Time 496204.96 μs
Device Time 3719352.30 μs
Self CPU Time 163763.41 μs
Self Device Time 3719352.30 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
sm80_xmma_fprop_implicit_gemm_indexed_f32f32_f32f32_f32_nchwkcrs_nchw_tilesize32x32x8_stage3_warpsize1x2x1_g1_ffma_aligna4_alignc4_execute_kernel__5x_cudnn
CPU Time 0.00 μs
Device Time 3719350.63 μs
Self CPU Time 0.00 μs
Self Device Time 3719350.63 μ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 4170413.05 μs
Device Time 72015.08 μs
Self CPU Time 4170413.05 μs
Self Device Time 72015.08 μ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 759152.28 μs
Device Time 483306.50 μs
Self CPU Time 14121.01 μ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 745033.59 μs
Device Time 483306.50 μs
Self CPU Time 22098.38 μs
Self Device Time 483306.50 μ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 
45291 warnings generated when compiling for host.
Suppressed 45325 warnings (45278 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_2/task_48/b3_s1_unroll_optimization/base/base.cu:29:5 bugprone-easily-swappable-parameters
29 | const float* __restrict__ output,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
30 | const float* __restrict__ scaling_factor,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
31 | const float* __restrict__ bias,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:29:31: note: the first parameter in the range is 'output'
29 | const float* __restrict__ output,
| ^~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:31:31: note: the last parameter in the range is 'bias'
31 | const float* __restrict__ bias,
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:49:21: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
49 | const int idx = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:69:19: warning: the parameter 'x' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
69 | torch::Tensor x,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:70: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]
70 | torch::Tensor conv_weight,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:71:5: warning: 2 adjacent parameters of 'forward' of similar type ('torch::Tensor') are easily swapped by mistake [bugprone-easily-swappable-parameters]
71 | torch::Tensor conv_bias,
| ^~~~~~~~~~~~~~~~~~~~~~~~
72 | torch::Tensor scaling_factor,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:71:19: note: the first parameter in the range is 'conv_bias'
71 | torch::Tensor conv_bias,
| ^~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:72:19: note: the last parameter in the range is 'scaling_factor'
72 | torch::Tensor scaling_factor,
| ^~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:72:19: warning: the parameter 'scaling_factor' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
72 | torch::Tensor scaling_factor,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:73:19: warning: the parameter 'bias' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
73 | torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:77:28: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
77 | const int batch_size = conv_out.size(0);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:78:30: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
78 | const int out_channels = conv_out.size(1);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:79:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
79 | const int depth = conv_out.size(2);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:80:24: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
80 | const int height = conv_out.size(3);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:81:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
81 | const int width = conv_out.size(4);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:89:45: warning: performing an implicit widening conversion to type 'unsigned long' of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
89 | conv3d_unroll_kernel<<<blocks, threads, 2 * out_channels * sizeof(float)>>>(
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:89:45: note: make conversion explicit to silence this warning
4 | conv3d_unroll_kernel<<<blocks, threads, 2 * out_channels * sizeof(float)>>>(
| ^~~~~~~~~~~~~~~~
| static_cast<unsigned long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s1_unroll_optimization/base/base.cu:89:45: note: perform multiplication in a wider type
89 | conv3d_unroll_kernel<<<blocks, threads, 2 * out_channels * sizeof(float)>>>(
| ^
| static_cast<long>( )