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 = 64
out_channels = 128
depth, height, width = 24, 48, 48
kernel_size = 3
stride = 2
padding = 1
min_value = -1.0
divisor = 2.0

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

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

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

_cpp_src = r"""
#include <torch/extension.h>
#include <ATen/ATen.h>
#include <vector>

torch::Tensor bias_clamp_div_inplace_cuda(torch::Tensor x, torch::Tensor bias, double min_value, double divisor);

torch::Tensor conv_transpose3d_bias_clamp_div(torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
                                              int64_t stride, int64_t padding,
                                              double min_value, double divisor) {
    c10::optional<at::Tensor> no_bias = c10::nullopt;
    at::Tensor y = at::conv_transpose3d(x, weight, no_bias,
                                        {stride, stride, stride},
                                        {padding, padding, padding},
                                        {0, 0, 0}, 1,
                                        {1, 1, 1});
    return bias_clamp_div_inplace_cuda(y, bias, min_value, divisor);
}
"""

_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void bias_clamp_div_kernel(float* __restrict__ p, const float* __restrict__ b,
                                      int64_t spatial, int channels, float minv, float invd) {
    int nc = blockIdx.x;
    int c = nc % channels;
    float bv = b[c];
    int64_t base = (int64_t)nc * spatial;
    for (int64_t s = threadIdx.x; s < spatial; s += blockDim.x) {
        float v = p[base + s] + bv;
        v = fmaxf(v, minv);
        p[base + s] = v * invd;
    }
}

torch::Tensor bias_clamp_div_inplace_cuda(torch::Tensor x, torch::Tensor bias, double min_value, double divisor) {
    TORCH_CHECK(x.is_cuda() && bias.is_cuda(), "tensors must be CUDA");
    TORCH_CHECK(x.scalar_type() == at::kFloat && bias.scalar_type() == at::kFloat, "only float32 supported");
    TORCH_CHECK(x.is_contiguous() && bias.is_contiguous(), "tensors must be contiguous");
    int64_t n = x.size(0);
    int channels = (int)x.size(1);
    int64_t spatial = x.size(2) * x.size(3) * x.size(4);
    if (x.numel() == 0) return x;
    float minv = static_cast<float>(min_value);
    float invd = 1.0f / static_cast<float>(divisor);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    bias_clamp_div_kernel<<<(int)(n * (int64_t)channels), 256, 0, stream>>>(x.data_ptr<float>(), bias.data_ptr<float>(), spatial, channels, minv, invd);
    return x;
}
"""

_ext = load_inline(
    name="convt_bias_clamp_div_ext_v1",
    cpp_sources=[_cpp_src],
    cuda_sources=[_cuda_src],
    functions=["conv_transpose3d_bias_clamp_div"],
    extra_cuda_cflags=["-O3"],
    extra_cflags=["-O3"],
    verbose=False,
)

class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, min_value, divisor):
        super(ModelNew, 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
        self._stride_i = stride if isinstance(stride, int) else stride[0]
        self._padding_i = padding if isinstance(padding, int) else padding[0]

    def forward(self, x):
        ct = self.conv_transpose
        return _ext.conv_transpose3d_bias_clamp_div(x, ct.weight, ct.bias, int(self._stride_i), int(self._padding_i), float(self.min_value), float(self.divisor))