← Back to Leaderboard

The AI CUDA Engineer 👷

100_ConvTranspose3d_Clamp_Min_Dividevectorized_stride_loop_base

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


def module_fn(
    x: torch.Tensor,
    stride: int,
    padding: int,
    min_value: float,
    divisor: float,
    conv_transpose: torch.Tensor,
    conv_transpose_bias: torch.Tensor,
) -> torch.Tensor:
    """
    Applies a transposed 3D convolution, clamps output to min value, and divides by constant.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, depth, height, width)
        stride (int): Stride of the transposed convolution
        padding (int): Padding of the transposed convolution
        min_value (float): Minimum value for clamping
        divisor (float): Value to divide output by
        conv_transpose (torch.Tensor): Transposed convolution weight tensor
        conv_transpose_bias (torch.Tensor): Bias tensor for transposed convolution

    Returns:
        torch.Tensor: Output tensor after applying transposed convolution, clamping and division
    """
    x = F.conv_transpose3d(
        x, conv_transpose, bias=conv_transpose_bias, stride=stride, padding=padding
    )
    x = torch.clamp(x, min=min_value)
    x = x / divisor
    return x


class Model(nn.Module):
    """
    A model that performs a transposed 3D convolution, clamps the output to a minimum value,
    and then divides the result by a constant.
    """

    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        min_value,
        divisor,
    ):
        super(Model, self).__init__()
        conv_transpose = nn.ConvTranspose3d(
            in_channels, out_channels, kernel_size, stride, padding
        )
        self.conv_transpose_parameter = conv_transpose.weight
        self.conv_transpose_bias = conv_transpose.bias

    def forward(self, x, stride, padding, min_value, divisor, fn=module_fn):
        return fn(
            x,
            stride,
            padding,
            min_value,
            divisor,
            self.conv_transpose_parameter,
            self.conv_transpose_bias,
        )


batch_size = 16
in_channels = 32
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
stride = 2
padding = 1
min_value = -1.0
divisor = 2.0


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


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

class Model(nn.Module):
    """
    A model that performs a transposed 3D convolution, clamps the output to a minimum value, 
    and then divides the result by a constant.
    """
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, min_value, divisor):
        super(Model, self).__init__()
        self.conv_transpose = nn.ConvTranspose3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)
        self.min_value = min_value
        self.divisor = divisor

    def forward(self, x):
        x = self.conv_transpose(x)
        x = torch.clamp(x, min=self.min_value)
        x = x / self.divisor
        return x

batch_size = 16
in_channels = 32
out_channels = 16
depth, height, width = 16, 32, 32
kernel_size = 3
stride = 2
padding = 1
min_value = -1.0
divisor = 2.0

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

def get_init_inputs():
    return [in_channels, out_channels, kernel_size, stride, padding, min_value, divisor]

Kernel Information

Related Kernels (Level 2, Task 100 • 100_ConvTranspose3d_Clamp_Min_Divide)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 optimized_clamp_divide_base 0.57 1.13 0.84
🥇 stride_loop_implementation_edit_1 0.57 1.13 0.84
🥇 modular_device_funcs3_base 0.57 1.13 0.84
🥇 balanced_workload_distribution_base 0.57 1.13 0.84
🥇 memory_coalescing_optimization_base 0.57 1.13 0.84
6 unrolled_convtranspose3d_edit_1 0.57 1.13 0.84
6 minimize_warp_divergence_edit_base 0.57 1.13 0.84
6 branchless_clamp_divide_opt_edit_1 0.57 1.13 0.84
6 modular_device_funcs2_edit_1 0.57 1.13 0.84
10 modular_device_funcs2_base 0.57 1.13 0.84
10 optimized_thread_block_indexing_base 0.57 1.13 0.84
10 optimized_thread_block_indexing_edit_1 0.57 1.13 0.84
13 vectorized_ldg_kernel_128_base 0.58 1.12 0.83
13 vectorized_ldg_kernel_128_edit_1 0.58 1.12 0.83
13 vectorized_modular_kernel_edit_1 0.58 1.12 0.83
13 vectorized_stride_loop_edit_1 0.58 1.12 0.83
13 vectorized_stride_loop_base 0.58 1.12 0.83
13 vectorized_modular_kernel_base 0.58 1.12 0.83
19 atomic_free_stride_loop_edit_1 0.58 1.12 0.83
19 modular_conv3d_clamp_divide_base 0.58 1.12 0.83
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>

template <typename scalar_t>
__global__ void clamp_and_divide_kernel(
    scalar_t* __restrict__ output,
    const int64_t numel,
    const float min_value,
    const float divisor) {
    
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int stride = gridDim.x * blockDim.x;
    
    // Vector processing for aligned portions
    float4* output4 = reinterpret_cast<float4*>(output);
    const int vec_elements = 4;
    const int vec_numel = numel / vec_elements;
    
    // Vector processing
    for (int idx = tid; idx < vec_numel; idx += stride) {
        float4 val4 = output4[idx];
        
        // Process each component
        val4.x = max(val4.x, min_value) / divisor;
        val4.y = max(val4.y, min_value) / divisor;
        val4.z = max(val4.z, min_value) / divisor;
        val4.w = max(val4.w, min_value) / divisor;
        
        output4[idx] = val4;
    }
    
    // Handle remaining elements
    const int remaining_start = vec_numel * vec_elements;
    for (int idx = remaining_start + tid; idx < numel; idx += stride) {
        output[idx] = max(output[idx], static_cast<scalar_t>(min_value)) / static_cast<scalar_t>(divisor);
    }
}

torch::Tensor forward(
    torch::Tensor input,
    int stride,
    int padding,
    float min_value,
    float divisor,
    torch::Tensor weight,
    torch::Tensor bias) {
    
    auto output = torch::conv_transpose3d(
        input,
        weight,
        bias,
        stride,
        padding
    );

    const int threads = 256;
    const int max_blocks = 65535;
    const int blocks = min(max_blocks, static_cast<int>((output.numel() + threads - 1) / threads));

    AT_DISPATCH_FLOATING_TYPES(output.scalar_type(), "clamp_and_divide", ([&] {
        clamp_and_divide_kernel<scalar_t><<<blocks, threads>>>(
            output.data_ptr<scalar_t>(),
            output.numel(),
            min_value,
            divisor
        );
    }));

    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward, "3D Transposed convolution with clamp and divide (CUDA)");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 1.228 inst/cycle 0.000 5
Executed Ipc Elapsed 1.184 inst/cycle 0.000 5
Issue Slots Busy 30.792 % 0.022 5
Issued Ipc Active 1.232 inst/cycle 0.000 5
SM Busy 30.792 % 0.022 5
Memory Throughput 2314794736629.820 byte/second 196340773812064813056.000 5
Mem Busy 38.588 % 0.044 5
Max Bandwidth 69.086 % 0.174 5
L1/TEX Hit Rate 45.628 % 0.004 5
L2 Hit Rate 50.408 % 0.011 5
Mem Pipes Busy 22.558 % 0.027 5
Warp Cycles Per Issued Instruction 35.372 cycle 0.034 5
Warp Cycles Per Executed Instruction 35.410 cycle 0.034 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 29.760 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 68.416 % 0.028 5
Achieved Active Warps Per SM 43.786 warp 0.011 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (20.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.
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 (68.4%) 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::conv_transpose3d
CPU Time 2494463.15 μs
Device Time 3516920.90 μs
Self CPU Time 12317.80 μ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 2482145.35 μs
Device Time 3516920.90 μs
Self CPU Time 17418.04 μ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 2464727.31 μs
Device Time 3516920.90 μs
Self CPU Time 36246.65 μ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_transpose
CPU Time 876422.31 μs
Device Time 2700886.28 μs
Self CPU Time 188257.21 μs
Self Device Time 2700886.28 μ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 3506941.12 μs
Device Time 61567.00 μs
Self CPU Time 3506941.12 μs
Self Device Time 61567.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
sm90_xmma_dgrad_implicit_gemm_indexed_f32f32_tf32f32_f32_nhwckrsc_nhwc_tilesize256x64x32_warpgroupsize1x1x1_g1_strided_execute_kernel__5x_cudnn
CPU Time 0.00 μs
Device Time 1618627.40 μs
Self CPU Time 0.00 μs
Self Device Time 1618627.40 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::add_
CPU Time 1541159.31 μs
Device Time 816034.62 μs
Self CPU Time 29317.71 μs
Self Device Time 816034.62 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
Status: Failed
45255 warnings and 1 error generated when compiling for host.
Error while processing /home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu.
Suppressed 45292 warnings (45245 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.
Found compiler error(s).
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:8:5 bugprone-easily-swappable-parameters
8 | const int64_t numel,
| ^~~~~~~~~~~~~~~~~~~~
9 | const float min_value,
| ~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:8:19: note: the first parameter in the range is 'numel'
8 | const int64_t numel,
| ^~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:9:17: note: the last parameter in the range is 'min_value'
9 | const float min_value,
| ^~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:8:5: note:
8 | const int64_t numel,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:9:5: note: 'const int64_t' and 'const float' may be implicitly converted: 'const int64_t' (as 'long') -> 'const float' (as 'float'), 'const float' (as 'float') -> 'const int64_t' (as 'long')
9 | const float min_value,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:12:21: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
12 | const int tid = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:13:24: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
13 | const int stride = gridDim.x * blockDim.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:18:27: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
18 | const int vec_numel = numel / vec_elements;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:41: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]
41 | torch::Tensor input,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:43:5: warning: 2 adjacent parameters of 'forward' of convertible types are easily swapped by mistake [bugprone-easily-swappable-parameters]
43 | int padding,
| ^~~~~~~~~~~~
44 | float min_value,
| ~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:43:9: note: the first parameter in the range is 'padding'
43 | int padding,
| ^~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:44:11: note: the last parameter in the range is 'min_value'
44 | float min_value,
| ^~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:44:5: note: 'int' and 'float' may be implicitly converted
44 | float min_value,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:46:19: warning: the parameter 'weight' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
46 | torch::Tensor weight,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:59:24: error: no matching function for call to 'min' [clang-diagnostic-error]
59 | const int blocks = min(max_blocks, static_cast<int>((output.numel() + threads - 1) / threads));
| ^~~
/home/common_modules/clang-tidy/20.0.0git/lib/clang/20/include/__clang_cuda_math.h:201:16: note: candidate function not viable: call to __device__ function from __host__ function
201 | __DEVICE__ int min(int __a, int __b) { return __nv_min(__a, __b); }
| ^
/usr/local/cuda/include/crt/math_functions.hpp:868:38: note: candidate function not viable: call to __device__ function from __host__ function
868 | __MATH_FUNCTIONS_DECL__ unsigned int min(const unsigned int a, const unsigned int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:873:38: note: candidate function not viable: call to __device__ function from __host__ function
873 | __MATH_FUNCTIONS_DECL__ unsigned int min(const int a, const unsigned int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:878:38: note: candidate function not viable: call to __device__ function from __host__ function
878 | __MATH_FUNCTIONS_DECL__ unsigned int min(const unsigned int a, const int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:883:34: note: candidate function not viable: call to __device__ function from __host__ function
883 | __MATH_FUNCTIONS_DECL__ long int min(const long int a, const long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:902:43: note: candidate function not viable: call to __device__ function from __host__ function
902 | __MATH_FUNCTIONS_DECL__ unsigned long int min(const unsigned long int a, const unsigned long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:919:43: note: candidate function not viable: call to __device__ function from __host__ function
919 | __MATH_FUNCTIONS_DECL__ unsigned long int min(const long int a, const unsigned long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:936:43: note: candidate function not viable: call to __device__ function from __host__ function
936 | __MATH_FUNCTIONS_DECL__ unsigned long int min(const unsigned long int a, const long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:953:39: note: candidate function not viable: call to __device__ function from __host__ function
953 | __MATH_FUNCTIONS_DECL__ long long int min(const long long int a, const long long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:958:48: note: candidate function not viable: call to __device__ function from __host__ function
958 | __MATH_FUNCTIONS_DECL__ unsigned long long int min(const unsigned long long int a, const unsigned long long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:963:48: note: candidate function not viable: call to __device__ function from __host__ function
963 | __MATH_FUNCTIONS_DECL__ unsigned long long int min(const long long int a, const unsigned long long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:968:48: note: candidate function not viable: call to __device__ function from __host__ function
968 | __MATH_FUNCTIONS_DECL__ unsigned long long int min(const unsigned long long int a, const long long int b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:973:31: note: candidate function not viable: call to __device__ function from __host__ function
973 | __MATH_FUNCTIONS_DECL__ float min(const float a, const float b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:978:32: note: candidate function not viable: call to __device__ function from __host__ function
978 | __MATH_FUNCTIONS_DECL__ double min(const double a, const double b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:983:32: note: candidate function not viable: call to __device__ function from __host__ function
983 | __MATH_FUNCTIONS_DECL__ double min(const float a, const double b)
| ^
/usr/local/cuda/include/crt/math_functions.hpp:988:32: note: candidate function not viable: call to __device__ function from __host__ function
988 | __MATH_FUNCTIONS_DECL__ double min(const double a, const float b)
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250211_optimize_b5_s4_e1_v2/level_2/task_100/b3_s1_vectorized_stride_loop/base/base.cu:61: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]
61 | AT_DISPATCH_FLOATING_TYPES(output.scalar_type(), "clamp_and_divide", ([&] {
| ^
/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__, \
| ^