```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

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

# Fused CUDA kernel: out = tanh(x - bias) with a per-channel bias over an NCHW tensor.
# This replaces the two separate elementwise PyTorch operators (sub + tanh) with a
# single memory pass, while the heavy transposed convolution stays on cuDNN.
sub_bias_tanh_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void sub_bias_tanh_kernel(const float* __restrict__ x,
                                     const float* __restrict__ bias,
                                     float* __restrict__ out,
                                     const int HW,
                                     const int C,
                                     const int num_planes) {
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    // Each blockIdx.y handles one (batch, channel) plane so no expensive
    // integer division is needed to recover the channel index.
    for (int plane = blockIdx.y; plane < num_planes; plane += gridDim.y) {
        const int c = plane % C;
        const long base = (long)plane * HW;
        if (i < HW) {
            out[base + i] = tanhf(x[base + i] - bias[c]);
        }
    }
}

torch::Tensor sub_bias_tanh_cuda(torch::Tensor x, torch::Tensor bias) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(bias.is_cuda(), "bias must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(bias.scalar_type() == at::kFloat, "bias must be float32");

    auto x_c = x.contiguous();
    auto bias_c = bias.contiguous();
    auto out = torch::empty_like(x_c);

    const int C  = (int)x_c.size(1);
    const int H  = (int)x_c.size(2);
    const int W  = (int)x_c.size(3);
    const int HW = H * W;
    const int num_planes = (int)x_c.size(0) * C;

    const int threads = 256;
    dim3 blocks((HW + threads - 1) / threads,
                num_planes < 65535 ? num_planes : 65535);

    sub_bias_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
        x_c.data_ptr<float>(),
        bias_c.data_ptr<float>(),
        out.data_ptr<float>(),
        HW, C, num_planes);

    return out;
}
"""

sub_bias_tanh_cpp_source = (
    "torch::Tensor sub_bias_tanh_cuda(torch::Tensor x, torch::Tensor bias);"
)

# Compile the fused (sub + tanh) CUDA operator
sub_bias_tanh = load_inline(
    name="sub_bias_tanh_fused",
    cpp_sources=sub_bias_tanh_cpp_source,
    cuda_sources=sub_bias_tanh_cuda_source,
    functions=["sub_bias_tanh_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the cuDNN-accelerated transposed convolution and
    replaces the bias subtraction + tanh activation with a single fused custom
    CUDA kernel (one memory pass instead of two).
    """
    def __init__(self, in_channels, out_channels, kernel_size, bias_shape, stride=2, padding=1, output_padding=1):
        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.sub_bias_tanh = sub_bias_tanh

    def forward(self, x):
        x = self.conv_transpose(x)
        x = self.sub_bias_tanh.sub_bias_tanh_cuda(x, self.bias)
        return x


batch_size = 32
in_channels  = 64
out_channels = 64
height = width = 256
kernel_size = 4
bias_shape = (out_channels, 1, 1)


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


def get_init_inputs():
    return [in_channels, out_channels, kernel_size, bias_shape]
```