```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
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel: out = tanh(x - bias) in a single memory pass.
# The ConvTranspose2d itself is left to cuDNN (implicit GEMM is hard to beat),
# while the two elementwise ops (bias-subtract + tanh) are fused to eliminate
# one full read+write pass over the large activation tensor.
fused_bias_tanh_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void bias_sub_tanh_kernel(
    const float* __restrict__ x,
    const float* __restrict__ bias,
    float* __restrict__ out,
    const int HW,
    const int C)
{
    // grid: (spatial_blocks, C, N); each thread handles 4 elements,
    // strided by blockDim.x so every load/store stays fully coalesced.
    const int base = blockIdx.x * (blockDim.x * 4) + threadIdx.x;
    const int c = blockIdx.y;
    const size_t slice = ((size_t)blockIdx.z * (size_t)C + (size_t)c) * (size_t)HW;
    const float b = bias[c];

    #pragma unroll
    for (int j = 0; j < 4; ++j) {
        const int hw = base + j * blockDim.x;
        if (hw < HW) {
            const size_t idx = slice + (size_t)hw;
            out[idx] = tanhf(x[idx] - b);
        }
    }
}

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

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

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

    const int threads = 256;
    const dim3 blocks((HW + threads * 4 - 1) / (threads * 4), C, N);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    bias_sub_tanh_kernel<<<blocks, threads, 0, stream>>>(
        x_c.data_ptr<float>(),
        b_c.data_ptr<float>(),
        out.data_ptr<float>(),
        HW, C);

    return out;
}
"""

fused_bias_tanh_cpp_source = (
    "torch::Tensor bias_sub_tanh_cuda(torch::Tensor x, torch::Tensor bias);"
)

fused_bias_tanh = load_inline(
    name="fused_bias_sub_tanh",
    cpp_sources=fused_bias_tanh_cpp_source,
    cuda_sources=fused_bias_tanh_cuda_source,
    functions=["bias_sub_tanh_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: ConvTranspose2d (cuDNN) + fused (bias subtract + tanh)
    custom CUDA kernel.
    """
    def __init__(self, in_channels, out_channels, kernel_size, bias_shape, stride=2, padding=1, output_padding=1):
        super(ModelNew, self).__init__()
        torch.backends.cudnn.benchmark = True
        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.fused_op = fused_bias_tanh

    def forward(self, x):
        x = self.conv_transpose(x)
        if x.is_cuda and x.dtype == torch.float32:
            # Single fused pass: tanh(x - bias)
            x = self.fused_op.bias_sub_tanh_cuda(x, self.bias)
        else:
            x = torch.tanh(x - self.bias)
        return x
```