← Back to Leaderboard

The AI CUDA Engineer 👷

2_ConvTranspose2d_BiasAdd_Clamp_Scaling_Clamp_Divideldg_128bit_align_opt_base

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


def module_fn(
    x: torch.Tensor,
    stride: int,
    padding: int,
    output_padding: int,
    scaling_factor: float,
    conv_transpose: torch.Tensor,
    conv_transpose_bias: torch.Tensor,
    bias: torch.Tensor,
) -> torch.Tensor:
    """Applies transposed convolution, bias addition, clamping, scaling, clamping and division.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, height, width)
        stride (int): Stride of the convolution
        padding (int): Zero-padding added to both sides of input
        output_padding (int): Additional size added to output shape
        scaling_factor (float): Factor to scale the tensor by
        conv_transpose (torch.Tensor): Transposed convolution weights
        conv_transpose_bias (torch.Tensor): Bias tensor for transposed convolution
        bias (torch.Tensor): Bias tensor to add after convolution

    Returns:
        torch.Tensor: Output tensor after applying operations
    """
    x = F.conv_transpose2d(
        x,
        conv_transpose,
        bias=conv_transpose_bias,
        stride=stride,
        padding=padding,
        output_padding=output_padding,
    )
    x = x + bias
    x = torch.clamp(x, min=0.0, max=1.0)
    x = x * scaling_factor
    x = torch.clamp(x, min=0.0, max=1.0)
    x = x / scaling_factor
    return x


class Model(nn.Module):
    """
    Model that performs a transposed convolution, adds a bias term, clamps, scales, clamps, and divides.
    """

    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        output_padding,
        bias_shape,
        scaling_factor,
    ):
        super(Model, self).__init__()
        conv_transpose = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            padding=padding,
            output_padding=output_padding,
        )
        self.conv_transpose_parameter = nn.Parameter(conv_transpose.weight)
        self.conv_tranpose_bias = nn.Parameter(conv_transpose.bias)
        self.bias_parameter = nn.Parameter(torch.randn(bias_shape) * 0.02)

    def forward(self, x, stride, padding, output_padding, scaling_factor, fn=module_fn):
        return fn(
            x,
            stride,
            padding,
            output_padding,
            scaling_factor,
            self.conv_transpose_parameter,
            self.conv_tranpose_bias,
            self.bias_parameter,
        )


batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3
stride = 2
padding = 1
output_padding = 1
bias_shape = (out_channels, 1, 1)
scaling_factor = 2.0


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


def get_init_inputs():
    return [
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        output_padding,
        bias_shape,
        scaling_factor,
    ]
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Model that performs a transposed convolution, adds a bias term, clamps, scales, clamps, and divides.
    """

    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        output_padding,
        bias_shape,
        scaling_factor,
    ):
        super(Model, self).__init__()
        self.conv_transpose = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
        )
        self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)

        self.scaling_factor = scaling_factor

    def forward(self, x):
        x = self.conv_transpose(x)
        x = x + self.bias
        x = torch.clamp(x, min=0.0, max=1.0)
        x = x * self.scaling_factor
        x = torch.clamp(x, min=0.0, max=1.0)
        x = x / self.scaling_factor
        return x


batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3
stride = 2
padding = 1
output_padding = 1
bias_shape = (out_channels, 1, 1)
scaling_factor = 2.0


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


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

Kernel Information

Related Kernels (Level 2, Task 2 • 2_ConvTranspose2d_BiasAdd_Clamp_Scaling_Clamp_Divide)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 shared_mem_sync_opt_base 0.18 1.63 0.77
🥇 ldg_128bit_align_opt_base 0.18 1.63 0.77
🥉 vec_ldg_128_align_base 0.18 1.62 0.77
🥉 combined_opt_base 0.18 1.62 0.77
🥉 modular_device_functions_opt_base 0.18 1.62 0.77
6 optimized_post_process_edit_1 0.18 1.62 0.76
6 modular_optimized_kernel_edit_1 0.18 1.62 0.76
6 coalesced_tiling_unrolled_kernel_base 0.18 1.62 0.76
6 optimized_post_process_base 0.18 1.62 0.76
6 aligned_ldg_128_opt_base 0.18 1.62 0.76
6 optimized_post_process_kernel_edit_1 0.18 1.62 0.76
6 optimized_post_process_kernel_base 0.18 1.62 0.76
6 modular_optimized_kernel_base 0.18 1.62 0.76
14 improved_coalesced_base 0.19 1.61 0.76
14 combined_post_process_base 0.19 1.61 0.76
14 coalesced_tiling_unrolled_kernel_edit_1 0.19 1.61 0.76
17 coalesced_tiling_kernel_edit_1 0.19 1.60 0.75
17 modular_coalesced_tiling_kernel_base 0.19 1.60 0.75
17 memory_coalescing_optimized_post_process_base 0.19 1.60 0.75
17 streamed_memory_transfer_2_convtranspose2d_biasadd_clamp_scaling_clamp_divide_edit_1 0.19 1.60 0.75
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>
#include <stdexcept>

#define BLOCK_SIZE 256

// Kernel that optimizes global memory load and store operations by using __ldg() for read-only accesses
// and aligning memory accesses to 128-bit boundaries using float4.
__global__ void ldg_128bit_align_kernel(
    float* __restrict__ output,
    const int total_size,    // total number of floats in the output
    const int height,
    const int width,
    const int channels,
    const float scaling_factor,
    const float* __restrict__ global_bias
) {
    int vec_size = total_size / 4;         // number of float4 chunks
    int remainder = total_size % 4;        // leftover elements

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    float4* out_vec = reinterpret_cast<float4*>(output);
    int hw_size = height * width;

    // Process vectorized elements in groups of 4 using a grid-stride loop
    for (int i = idx; i < vec_size; i += stride) {
        // Load 4 floats at once using __ldg for enhanced read-only performance
        float4 data = __ldg(&out_vec[i]);
        int base_index = i * 4;
        float results[4];
        
        // Process each element of the float4
        {
            int index = base_index;
            int c = (index / hw_size) % channels;
            float val = data.x + __ldg(&global_bias[c]);
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            val = val * scaling_factor;
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            results[0] = val / scaling_factor;
        }
        {
            int index = base_index + 1;
            int c = (index / hw_size) % channels;
            float val = data.y + __ldg(&global_bias[c]);
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            val = val * scaling_factor;
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            results[1] = val / scaling_factor;
        }
        {
            int index = base_index + 2;
            int c = (index / hw_size) % channels;
            float val = data.z + __ldg(&global_bias[c]);
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            val = val * scaling_factor;
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            results[2] = val / scaling_factor;
        }
        {
            int index = base_index + 3;
            int c = (index / hw_size) % channels;
            float val = data.w + __ldg(&global_bias[c]);
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            val = val * scaling_factor;
            val = fminf(fmaxf(val, 0.0f), 1.0f);
            results[3] = val / scaling_factor;
        }

        // Write the processed values back in a vectorized manner
        float4 out_val = make_float4(results[0], results[1], results[2], results[3]);
        out_vec[i] = out_val;
    }

    // Process any remaining elements that weren't covered by the vectorized loop
    int rem_start = vec_size * 4;
    for (int i = idx; i < remainder; i += stride) {
        int index = rem_start + i;
        int c = (index / hw_size) % channels;
        float val = __ldg(&output[index]) + __ldg(&global_bias[c]);
        val = fminf(fmaxf(val, 0.0f), 1.0f);
        val = val * scaling_factor;
        val = fminf(fmaxf(val, 0.0f), 1.0f);
        output[index] = val / scaling_factor;
    }
}

// Forward function performs conv_transpose2d followed by the post-processing kernel

torch::Tensor forward(
    torch::Tensor x,
    int64_t stride,
    int64_t padding,
    int64_t output_padding,
    float scaling_factor,
    torch::Tensor conv_transpose,
    torch::Tensor conv_transpose_bias,
    torch::Tensor bias
) {
    // Perform transposed convolution using PyTorch's built-in function
    auto output = torch::conv_transpose2d(
        x, conv_transpose, conv_transpose_bias,
        stride, padding, output_padding
    );

    const int batch_size = output.size(0);
    const int channels = output.size(1);
    const int height = output.size(2);
    const int width = output.size(3);
    const int total_size = batch_size * channels * height * width;

    int threads = BLOCK_SIZE;
    int vec_size = total_size / 4;  // Number of float4 groups
    int blocks = (vec_size + threads - 1) / threads;
    if (blocks == 0) blocks = 1;

    ldg_128bit_align_kernel<<<blocks, threads>>>(
        output.data_ptr<float>(),
        total_size,
        height,
        width,
        channels,
        scaling_factor,
        bias.data_ptr<float>()
    );

    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward, "Optimized post-processing kernel with __ldg() and 128-bit alignment (CUDA)");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 2.760 inst/cycle 0.000 5
Executed Ipc Elapsed 2.462 inst/cycle 0.001 5
Issue Slots Busy 69.212 % 0.007 5
Issued Ipc Active 2.766 inst/cycle 0.000 5
SM Busy 69.212 % 0.007 5
Memory Throughput 1842253282235.698 byte/second 679779132915798507520.000 5
Mem Busy 38.156 % 0.301 5
Max Bandwidth 54.998 % 0.605 5
L1/TEX Hit Rate 55.244 % 0.000 5
L2 Hit Rate 51.306 % 0.016 5
Mem Pipes Busy 15.748 % 0.051 5
Warp Cycles Per Issued Instruction 19.386 cycle 0.000 5
Warp Cycles Per Executed Instruction 19.444 cycle 0.000 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 26.990 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 84.376 % 0.004 5
Achieved Active Warps Per SM 54.004 warp 0.002 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (56.6%) 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 (84.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::conv_transpose2d
CPU Time 2693340.66 μs
Device Time 2318993.98 μs
Self CPU Time 20894.62 μ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 2672446.04 μs
Device Time 2318993.98 μs
Self CPU Time 31099.54 μ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 2641346.50 μs
Device Time 2318993.98 μs
Self CPU Time 66409.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 2269303.25 μs
Device Time 1872447.79 μs
Self CPU Time 448090.10 μs
Self Device Time 1872447.79 μ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 1663616.19 μs
Device Time 25624.02 μs
Self CPU Time 1663616.19 μs
Self Device Time 25624.02 μ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 122550.34 μs
Device Time 1142479.80 μs
Self CPU Time 25792.42 μ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
Status: Completed
45294 warnings generated when compiling for host.
Suppressed 45328 warnings (45281 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_2/b10_s1_ldg_128bit_align_opt/base/base.cu:13:5 bugprone-easily-swappable-parameters
13 | const int total_size, // total number of floats in the output
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
14 | const int height,
| ~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:13:15: note: the first parameter in the range is 'total_size'
13 | const int total_size, // total number of floats in the output
| ^~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:14:15: note: the last parameter in the range is 'height'
14 | const int height,
| ^~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:15:5: warning: 3 adjacent parameters of 'ldg_128bit_align_kernel' of convertible types are easily swapped by mistake [bugprone-easily-swappable-parameters]
15 | const int width,
| ^~~~~~~~~~~~~~~~
16 | const int channels,
| ~~~~~~~~~~~~~~~~~~~
17 | const float scaling_factor,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:15:15: note: the first parameter in the range is 'width'
15 | const int width,
| ^~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:17:17: note: the last parameter in the range is 'scaling_factor'
17 | const float scaling_factor,
| ^~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:17:5: note: 'const int' and 'const float' may be implicitly converted: 'const int' (as 'int') -> 'const float' (as 'float'), 'const float' (as 'float') -> 'const int' (as 'int')
17 | const float scaling_factor,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:23:15: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
23 | int idx = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:24:18: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
24 | int stride = blockDim.x * gridDim.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:95: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]
95 | torch::Tensor x,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:98:5: warning: 2 adjacent parameters of 'forward' of convertible types are easily swapped by mistake [bugprone-easily-swappable-parameters]
98 | int64_t output_padding,
| ^~~~~~~~~~~~~~~~~~~~~~~
99 | float scaling_factor,
| ~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:98:13: note: the first parameter in the range is 'output_padding'
98 | int64_t output_padding,
| ^~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:99:11: note: the last parameter in the range is 'scaling_factor'
99 | float scaling_factor,
| ^~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:98:5: note:
98 | int64_t output_padding,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:99:5: note: 'int64_t' and 'float' may be implicitly converted: 'int64_t' (as 'long') -> 'float', 'float' -> 'int64_t' (as 'long')
99 | float scaling_factor,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:100:19: warning: the parameter 'conv_transpose' 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_transpose,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:101:5: warning: 2 adjacent parameters of 'forward' of similar type ('torch::Tensor') are easily swapped by mistake [bugprone-easily-swappable-parameters]
101 | torch::Tensor conv_transpose_bias,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
102 | torch::Tensor bias
| ~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:101:19: note: the first parameter in the range is 'conv_transpose_bias'
101 | torch::Tensor conv_transpose_bias,
| ^~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:102:19: note: the last parameter in the range is 'bias'
102 | torch::Tensor bias
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:102: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]
102 | torch::Tensor bias
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:110:28: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
110 | const int batch_size = output.size(0);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:111:26: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
111 | const int channels = output.size(1);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:112:24: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
112 | const int height = output.size(2);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_2/b10_s1_ldg_128bit_align_opt/base/base.cu:113:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
113 | const int width = output.size(3);
| ^