← Back to Leaderboard

The AI CUDA Engineer 👷

48_Conv3d_Scaling_Tanh_Multiply_Sigmoidconstant_memory_optimization_edit_1

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]

Kernel Information

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>

// Constant memory for scaling factors and biases
__constant__ float const_scaling_factor[1024];
__constant__ float const_bias[1024];

// CUDA kernel that uses constant memory for frequently accessed parameters
__global__ void conv3d_constant_memory_kernel(
    const float* __restrict__ output,
    float* __restrict__ result,
    const int batch_size,
    const int out_channels,
    const int depth,
    const int height, 
    const int width) {

    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 = __ldg(&output[idx]);
        val *= const_scaling_factor[c];
        val = tanh(val);
        val *= const_bias[c];
        val = 1.0f / (1.0f + exp(-val));

        result[idx] = val;
    }
}

// Forward function that performs conv3d followed by scaling, tanh, bias multiplication and sigmoid
// It launches the optimized kernel using constant memory

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);

    // Copy scaling factors and biases to constant memory
    cudaMemcpyToSymbol(const_scaling_factor, scaling_factor.data_ptr<float>(), out_channels * sizeof(float));
    cudaMemcpyToSymbol(const_bias, bias.data_ptr<float>(), out_channels * sizeof(float));

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

    conv3d_constant_memory_kernel<<<blocks, threads>>>(
        conv_out.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, "Optimized Conv3d scale tanh bias sigmoid forward using constant memory");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 3.250 inst/cycle 0.000 5
Executed Ipc Elapsed 3.178 inst/cycle 0.000 5
Issue Slots Busy 81.262 % 0.001 5
Issued Ipc Active 3.250 inst/cycle 0.000 5
SM Busy 81.262 % 0.001 5
Memory Throughput 1645593389874.222 byte/second 13248074707051128832.000 5
Mem Busy 27.256 % 0.003 5
Max Bandwidth 49.104 % 0.011 5
L1/TEX Hit Rate 0.000 % 0.000 5
L2 Hit Rate 49.894 % 0.001 5
Mem Pipes Busy 53.672 % 0.009 5
Warp Cycles Per Issued Instruction 15.170 cycle 0.001 5
Warp Cycles Per Executed Instruction 15.174 cycle 0.001 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 28.650 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 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 78.028 % 0.000 5
Achieved Active Warps Per SM 49.940 warp 0.000 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (56.5%) 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 (78.1%) 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 550345.29 μs
Device Time 2969.14 μs
Self CPU Time 63.41 μ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 567661.25 μs
Device Time 0.00 μs
Self CPU Time 20637.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::conv3d
CPU Time 517778.92 μs
Device Time 4315107.04 μs
Self CPU Time 13288.60 μ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 504490.32 μs
Device Time 4315107.04 μs
Self CPU Time 16543.82 μ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 487946.50 μs
Device Time 4315107.04 μs
Self CPU Time 33108.66 μ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 381988.49 μs
Device Time 3745026.81 μs
Self CPU Time 322070.61 μs
Self Device Time 3745026.81 μ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 3745024.92 μs
Self CPU Time 0.00 μs
Self Device Time 3745024.92 μ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 4052824.72 μs
Device Time 70358.05 μs
Self CPU Time 4052824.72 μs
Self Device Time 70358.05 μ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 3588378.28 μs
Device Time 486948.17 μs
Self CPU Time 15024.71 μ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 3573355.69 μs
Device Time 486948.17 μs
Self CPU Time 21607.90 μs
Self Device Time 486948.17 μ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
45289 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_s0_constant_memory_optimization/edit_1/edit_1.cu:21:21 bugprone-narrowing-conversions
21 | 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_s0_constant_memory_optimization/edit_1/edit_1.cu:41: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]
41 | torch::Tensor x,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:42: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]
42 | torch::Tensor conv_weight,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:43:5: warning: 2 adjacent parameters of 'forward' of similar type ('torch::Tensor') are easily swapped by mistake [bugprone-easily-swappable-parameters]
43 | torch::Tensor conv_bias,
| ^~~~~~~~~~~~~~~~~~~~~~~~
44 | torch::Tensor scaling_factor,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:43:19: note: the first parameter in the range is 'conv_bias'
43 | torch::Tensor conv_bias,
| ^~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:44:19: note: the last parameter in the range is 'scaling_factor'
44 | torch::Tensor scaling_factor,
| ^~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:44: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]
44 | torch::Tensor scaling_factor,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:45: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]
45 | torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_48/b3_s0_constant_memory_optimization/edit_1/edit_1.cu:49:28: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
49 | 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_s0_constant_memory_optimization/edit_1/edit_1.cu:50:30: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
50 | 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_s0_constant_memory_optimization/edit_1/edit_1.cu:51:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
51 | 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_s0_constant_memory_optimization/edit_1/edit_1.cu:52:24: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
52 | 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_s0_constant_memory_optimization/edit_1/edit_1.cu:53:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
53 | const int width = conv_out.size(4);
| ^