import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Model that performs a transposed convolution, applies softmax, adds a bias term, scales the result, and applies sigmoid.
    """
    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)) 
        self.scaling_factor = scaling_factor

    def forward(self, x):
        x = self.conv_transpose(x)
        x = torch.softmax(x, dim=1)
        x = x + self.bias
        x = x * self.scaling_factor
        x = torch.sigmoid(x)
        return x

batch_size = 128
in_channels = 64
out_channels = 128
height, width = 64, 64
kernel_size = 4
stride = 2
padding = 1
output_padding = 1
bias_shape = (out_channels, 1, 1)
scaling_factor = 2.0

def get_inputs():
    return [torch.rand(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]

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_source = r'''
#include <torch/extension.h>
torch::Tensor fused_postprocess_cuda(torch::Tensor x, torch::Tensor bias, double scaling);
torch::Tensor fused_postprocess(torch::Tensor x, torch::Tensor bias, double scaling) { return fused_postprocess_cuda(x, bias, scaling); }
'''

_cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <float.h>

constexpr int PIX = 32;
constexpr int GRP = 8;
constexpr int COUT = 128;

__global__ void fused_post_kernel(const float* __restrict__ x,
                                  const float* __restrict__ bias,
                                  float* __restrict__ out,
                                  int H, int W, float scale) {
    __shared__ float partial[GRP][PIX];
    __shared__ float maxv[PIX];
    __shared__ float sumv[PIX];
    __shared__ float ebuf[COUT][PIX];
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int HW = H * W;
    int n = blockIdx.y;
    int s = blockIdx.x * PIX + tx;
    bool valid = s < HW;

    float m = -FLT_MAX;
    if (valid) {
        #pragma unroll
        for (int c = ty; c < COUT; c += GRP) {
            float v = x[((n * COUT + c) * HW) + s];
            m = fmaxf(m, v);
        }
    }
    partial[ty][tx] = m;
    __syncthreads();
    if (ty == 0) {
        float mm = partial[0][tx];
        #pragma unroll
        for (int g = 1; g < GRP; ++g) mm = fmaxf(mm, partial[g][tx]);
        maxv[tx] = mm;
    }
    __syncthreads();

    float ss = 0.0f;
    if (valid) {
        float mm = maxv[tx];
        #pragma unroll
        for (int c = ty; c < COUT; c += GRP) {
            float e = __expf(x[((n * COUT + c) * HW) + s] - mm);
            ebuf[c][tx] = e;
            ss += e;
        }
    }
    partial[ty][tx] = ss;
    __syncthreads();
    if (ty == 0) {
        float ssum = partial[0][tx];
        #pragma unroll
        for (int g = 1; g < GRP; ++g) ssum += partial[g][tx];
        sumv[tx] = ssum;
    }
    __syncthreads();

    if (valid) {
        float inv = 1.0f / sumv[tx];
        #pragma unroll
        for (int c = ty; c < COUT; c += GRP) {
            int idx = ((n * COUT + c) * HW) + s;
            float sm = ebuf[c][tx] * inv;
            float z = (sm + bias[c]) * scale;
            out[idx] = 1.0f / (1.0f + __expf(-z));
        }
    }
}

torch::Tensor fused_postprocess_cuda(torch::Tensor x, torch::Tensor bias, double scaling) {
    auto out = torch::empty_like(x);
    int N = (int)x.size(0);
    int H = (int)x.size(2);
    int W = (int)x.size(3);
    int HW = H * W;
    dim3 block(PIX, GRP);
    dim3 grid((HW + PIX - 1) / PIX, N);
    fused_post_kernel<<<grid, block>>>(x.data_ptr<float>(), bias.data_ptr<float>(), out.data_ptr<float>(), H, W, (float)scaling);
    return out;
}
'''

_ext = load_inline(name='fused_ct_post_ext_final', cpp_sources=[_cpp_source], cuda_sources=[_cuda_source], functions=['fused_postprocess'], verbose=False, extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])

class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, bias_shape, scaling_factor):
        super(ModelNew, 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))
        self.scaling_factor = scaling_factor
    def forward(self, x):
        y = self.conv_transpose(x)
        return _ext.fused_postprocess(y, self.bias, float(self.scaling_factor))