The AI CUDA Engineer 👷

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,
    conv_transpose: torch.Tensor,
    conv_transpose_bias: torch.Tensor,
    multiplier: float,
) -> torch.Tensor:
    """
    Applies transposed convolution, scalar multiplication, and multiple global average pooling operations.

    Args:
        x (torch.Tensor): Input tensor of shape (batch_size, in_channels, height, width)
        stride (int): Stride of the transposed convolution
        padding (int): Padding of the transposed convolution
        output_padding (int): Additional size added to output shape
        conv_transpose (torch.Tensor): Transposed convolution weight tensor
        conv_transpose_bias (torch.Tensor): Bias tensor for transposed convolution
        multiplier (float): Scalar multiplier value

    Returns:
        torch.Tensor: Scalar output after applying operations
    """
    x = F.conv_transpose2d(
        x,
        conv_transpose,
        bias=conv_transpose_bias,
        stride=stride,
        padding=padding,
        output_padding=output_padding,
    )
    x = x * multiplier
    x = torch.mean(x, dim=[2, 3], keepdim=True)
    x = torch.mean(x, dim=[2, 3], keepdim=True)
    x = torch.mean(x)
    return x


class Model(nn.Module):
    """
    Model that performs a transposed convolution, multiplies by a scalar, applies global average pooling,
    another global average pooling, and then calculates the mean.
    """

    def __init__(
        self,
        in_channels,
        out_channels,
        kernel_size,
        stride,
        padding,
        output_padding,
        multiplier,
    ):
        super(Model, self).__init__()
        conv = nn.ConvTranspose2d(
            in_channels,
            out_channels,
            kernel_size,
            stride=stride,
            padding=padding,
            output_padding=output_padding,
        )
        self.conv_transpose_parameter = nn.Parameter(conv.weight)
        self.conv_transpose_bias = nn.Parameter(
            conv.bias
            + torch.randn(
                conv.bias.shape, device=conv.bias.device, dtype=conv.bias.dtype
            )
            * 0.02
        )
        self.multiplier = multiplier

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


batch_size = 128
in_channels = 3
out_channels = 16
height, width = 32, 32
kernel_size = 3
stride = 2
padding = 1
output_padding = 1
multiplier = 0.5


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


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

#include <torch/extension.h>
#include <ATen/ATen.h>

__global__ void optimized_avg_pool_kernel(const float* __restrict__ input, float* __restrict__ output, int height, int width, int channels) {
    extern __shared__ float shared_data[];
    const unsigned int tid = threadIdx.x;
    const unsigned int bid = blockIdx.x;
    const unsigned int num_threads = blockDim.x;
    const unsigned int lane_id = tid % 32;
    const unsigned int warp_id = tid / 32;
    const unsigned int num_warps = num_threads / 32;
    
    // Each thread initializes its accumulator
    float sum = 0.0f;
    
    // Coalesced global memory loads with grid-stride loop
    const unsigned int total_elements = height * width;
    const unsigned int grid_stride = num_threads * gridDim.x;
    const unsigned int start_idx = bid * total_elements + tid;
    
    // Direct accumulation without intermediate shared memory
    for (unsigned int i = start_idx; i < (bid + 1) * total_elements; i += grid_stride) {
        sum += input[i];
    }
    
    // Warp reduction using shuffle
    #pragma unroll
    for (int offset = 16; offset > 0; offset /= 2) {
        sum += __shfl_down_sync(0xffffffff, sum, offset);
    }
    
    // First thread in each warp writes to shared memory
    if (lane_id == 0) {
        shared_data[warp_id] = sum;
    }
    __syncthreads();
    
    // Final reduction across warps
    if (tid < num_warps) {
        float warp_sum = shared_data[tid];
        #pragma unroll
        for (int offset = num_warps/2; offset > 0; offset /= 2) {
            warp_sum += __shfl_down_sync(0xffffffff, warp_sum, offset);
        }
        
        if (tid == 0) {
            output[bid] = warp_sum / (height * width);
        }
    }
}

at::Tensor module_fn(
    at::Tensor x,
    int64_t stride,
    int64_t padding,
    int64_t output_padding,
    at::Tensor conv_transpose,
    at::Tensor conv_transpose_bias,
    double multiplier
) {
    // Apply transposed convolution
    at::Tensor y = at::conv_transpose2d(
        x,
        conv_transpose,
        conv_transpose_bias,
        {stride, stride},             // stride
        {padding, padding},           // padding
        {output_padding, output_padding}, // output_padding
        1,                            // groups
        {1, 1}                        // dilation
    );

    // Element-wise multiplication by scalar multiplier
    y = y * multiplier;

    // Prepare output tensor for global average pooling
    auto output = at::zeros({y.size(0), y.size(1)}, y.options());

    // Launch kernel for optimized global average pooling
    int threads = 256;
    int blocks = y.size(0) * y.size(1);
    int shared_mem_size = threads * sizeof(float);
    optimized_avg_pool_kernel<<<blocks, threads, shared_mem_size>>>(y.data_ptr<float>(), output.data_ptr<float>(), y.size(2), y.size(3), y.size(1));

    // Compute mean over all elements to get scalar output
    auto result = output.mean();

    return result;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &module_fn, "Module function");
}