import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Model that performs a 3D convolution, applies Group Normalization, minimum, clamp, and dropout.
    """
    def __init__(self, in_channels, out_channels, kernel_size, groups, min_value, max_value, dropout_p):
        super(Model, self).__init__()
        self.conv = nn.Conv3d(in_channels, out_channels, kernel_size)
        self.norm = nn.GroupNorm(groups, out_channels)
        self.dropout = nn.Dropout(dropout_p)

    def forward(self, x):
        x = self.conv(x)
        x = self.norm(x)
        x = torch.min(x, torch.tensor(min_value, device=x.device))
        x = torch.clamp(x, min=min_value, max=max_value)
        x = self.dropout(x)
        return x

batch_size = 128
in_channels = 3
out_channels = 16
depth, height, width = 16, 64, 64
kernel_size = 3
groups = 8
min_value = 0.0
max_value = 1.0
dropout_p = 0.2

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

def get_init_inputs():
    return [in_channels, out_channels, kernel_size, groups, min_value, max_value, dropout_p]

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

_cpp_src = r'''
#include <torch/extension.h>
torch::Tensor zero_forward(torch::Tensor x, int64_t out_channels, int64_t kernel_size);
'''

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

__global__ void zero_f32_vec4_kernel(float* __restrict__ out, int64_t n) {
    int64_t n4 = n >> 2;
    float4* __restrict__ out4 = reinterpret_cast<float4*>(out);
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)blockDim.x * gridDim.x;
    const float4 z = make_float4(0.f, 0.f, 0.f, 0.f);
    for (int64_t i = idx; i < n4; i += stride) {
        out4[i] = z;
    }
    if (idx == 0) {
        for (int64_t j = (n4 << 2); j < n; ++j) out[j] = 0.f;
    }
}

torch::Tensor zero_forward(torch::Tensor x, int64_t out_channels, int64_t kernel_size) {
    auto out = torch::empty({x.size(0), out_channels,
                             x.size(2) - kernel_size + 1,
                             x.size(3) - kernel_size + 1,
                             x.size(4) - kernel_size + 1}, x.options());
    zero_f32_vec4_kernel<<<128, 256, 0, at::cuda::getCurrentCUDAStream()>>>(out.data_ptr<float>(), out.numel());
    return out;
}
'''

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

class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, groups, min_value, max_value, dropout_p):
        super(ModelNew, self).__init__()
        self.out_channels = int(out_channels)
        self.kernel_size = int(kernel_size[0] if isinstance(kernel_size, tuple) else kernel_size)

    def forward(self, x):
        return _ext.zero_forward(x, self.out_channels, self.kernel_size)