import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs a convolution, takes the minimum with a constant, adds a bias term, and multiplies by a scaling factor.
    """
    def __init__(self, in_channels, out_channels, kernel_size, constant_value, bias_shape, scaling_factor):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.constant_value = constant_value
        self.bias = nn.Parameter(torch.randn(bias_shape))
        self.scaling_factor = scaling_factor

    def forward(self, x):
        x = self.conv(x)
        x = torch.min(x, torch.tensor(self.constant_value))
        x = x + self.bias
        x = x * self.scaling_factor
        return x

batch_size = 128
in_channels = 64
out_channels = 128
height = width = 128
kernel_size = 3
constant_value = 0.5
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, constant_value, bias_shape, scaling_factor]

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

_cpp_src = r'''
#include <torch/extension.h>
torch::Tensor fused_postprocess(torch::Tensor x, torch::Tensor bias, double constant, double scale);
'''

_cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>

__global__ void fused_postprocess_vec4_plane_kernel(float* __restrict__ x,
                                                    const float* __restrict__ bias,
                                                    const int C,
                                                    const int HW4,
                                                    const float constant,
                                                    const float scale) {
    const int plane = blockIdx.x;
    const float b = bias[plane % C];
    float4* __restrict__ xp4 = reinterpret_cast<float4*>(x + ((long)plane * HW4 * 4));
    for (int j = threadIdx.x; j < HW4; j += blockDim.x) {
        float4 v = xp4[j];
        v.x = (fminf(v.x, constant) + b) * scale;
        v.y = (fminf(v.y, constant) + b) * scale;
        v.z = (fminf(v.z, constant) + b) * scale;
        v.w = (fminf(v.w, constant) + b) * scale;
        xp4[j] = v;
    }
}

__global__ void fused_postprocess_scalar_kernel(float* __restrict__ x,
                                                const float* __restrict__ bias,
                                                const long total,
                                                const int C,
                                                const int HW,
                                                const float constant,
                                                const float scale) {
    long idx = blockIdx.x * blockDim.x + threadIdx.x;
    const long stride = (long)blockDim.x * gridDim.x;
    for (; idx < total; idx += stride) {
        int c = (idx / HW) % C;
        float v = x[idx];
        x[idx] = (fminf(v, constant) + bias[c]) * scale;
    }
}

torch::Tensor fused_postprocess(torch::Tensor x, torch::Tensor bias, double constant, double scale) {
    const c10::cuda::CUDAGuard device_guard(x.device());
    const long total = x.numel();
    if (total == 0) return x;
    const int N = (int)x.size(0);
    const int C = (int)x.size(1);
    const int H = (int)x.size(2);
    const int W = (int)x.size(3);
    const int HW = H * W;
    const float cst = (float)constant;
    const float scl = (float)scale;
    const int threads = 256;
    if ((HW & 3) == 0 && (((uintptr_t)x.data_ptr<float>()) & 15) == 0) {
        const int HW4 = HW >> 2;
        fused_postprocess_vec4_plane_kernel<<<N * C, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(), C, HW4, cst, scl);
    } else {
        int blocks = (int)((total + threads - 1) / threads);
        if (blocks > 65535) blocks = 65535;
        fused_postprocess_scalar_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
            x.data_ptr<float>(), bias.data_ptr<float>(), total, C, HW, cst, scl);
    }
    return x;
}
'''

_ext = load_inline(name='fused_min_bias_scale_ext_v3',
                   cpp_sources=_cpp_src,
                   cuda_sources=_cuda_src,
                   functions=['fused_postprocess'],
                   extra_cuda_cflags=['-O3'],
                   extra_cflags=['-O3'],
                   verbose=False)

class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, constant_value, bias_shape, scaling_factor):
        super(ModelNew, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.constant_value = float(constant_value)
        self.bias = nn.Parameter(torch.randn(bias_shape))
        self.scaling_factor = float(scaling_factor)

    def forward(self, x):
        y = self.conv(x)
        return _ext.fused_postprocess(y, self.bias, self.constant_value, self.scaling_factor)