Looking at the profiler, the big costs beyond the GEMM itself are: a 3.2ms bias-add elementwise pass, a 3.1ms NHWC→NCHW output conversion, and my 3.1ms epilogue. I can eliminate the layout conversion by keeping everything channels_last (the elementwise epilogue doesn't care about logical layout), and fuse the conv bias into my epilogue kernel so ATen's separate bias pass disappears.

```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

torch.backends.cudnn.benchmark = True
torch.backends.cudnn.benchmark_limit = 0  # exhaustive cudnn algo search
torch.backends.cudnn.allow_tf32 = True
torch.backends.cuda.matmul.allow_tf32 = True

# Fused epilogue operating directly on the channels_last (NHWC) output of the
# transposed convolution. Since the epilogue is element-wise, the NHWC layout
# is irrelevant except for mapping a flat index -> channel for the bias add,
# which in NHWC is simply (flat_idx % C). This kernel fuses:
#   bias add -> Mish -> add_value -> Hardtanh -> scale
# eliminating ATen's separate bias pass and cudnn's NHWC->NCHW transpose.
fused_epilogue_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float epilogue(float v, float b, float add_value, float scale) {
    float x = v + b;
    // mish(x) = x * tanh(softplus(x)) = x * u / (u + 2), u = t*(t+2), t = exp(x)
    float t = expf(fminf(x, 20.0f));
    float u = t * (t + 2.0f);
    float m = x * (u / (u + 2.0f));
    float y = m + add_value;
    y = fminf(fmaxf(y, -1.0f), 1.0f);  // hardtanh in [-1, 1]
    return y * scale;
}

// Vectorized path (requires numel % 4 == 0 and C % 4 == 0): each float4 spans
// 4 consecutive channels in NHWC, so the bias is also read as an aligned float4.
__global__ void fused_epilogue_nhwc_vec4_kernel(const float4* __restrict__ in,
                                                float4* __restrict__ out,
                                                const float* __restrict__ bias,
                                                long n4, long c4,
                                                float add_value, float scale) {
    long stride = (long)gridDim.x * blockDim.x;
    for (long v = (long)blockIdx.x * blockDim.x + threadIdx.x; v < n4; v += stride) {
        long cbase = (v % c4) << 2;
        float4 b = *reinterpret_cast<const float4*>(bias + cbase);
        float4 val = in[v];
        float4 r;
        r.x = epilogue(val.x, b.x, add_value, scale);
        r.y = epilogue(val.y, b.y, add_value, scale);
        r.z = epilogue(val.z, b.z, add_value, scale);
        r.w = epilogue(val.w, b.w, add_value, scale);
        out[v] = r;
    }
}

// Generic fallback path.
__global__ void fused_epilogue_nhwc_kernel(const float* __restrict__ in,
                                           float* __restrict__ out,
                                           const float* __restrict__ bias,
                                           long n, long C,
                                           float add_value, float scale) {
    long stride = (long)gridDim.x * blockDim.x;
    for (long i = (long)blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride) {
        long c = i % C;
        out[i] = epilogue(in[i], bias[c], add_value, scale);
    }
}

torch::Tensor fused_epilogue_nhwc_cuda(torch::Tensor x, torch::Tensor bias, double add_value, double scale) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(x.is_contiguous(at::MemoryFormat::ChannelsLast), "x must be channels_last contiguous");
    TORCH_CHECK(bias.is_cuda() && bias.scalar_type() == torch::kFloat32, "bias must be CUDA float32");

    auto b = bias.contiguous();
    const long C = x.size(1);
    TORCH_CHECK(b.numel() == C, "bias size mismatch");

    auto out = torch::empty_like(x);  // preserves channels_last layout
    const long n = x.numel();
    if (n == 0) return out;

    auto stream = at::cuda::getCurrentCUDAStream();
    const int block = 256;

    if ((n % 4 == 0) && (C % 4 == 0)) {
        const long n4 = n >> 2;
        long blocks = (n4 + block - 1) / block;
        if (blocks > 2147483647L) blocks = 2147483647L;
        fused_epilogue_nhwc_vec4_kernel<<<(unsigned int)blocks, block, 0, stream>>>(
            reinterpret_cast<const float4*>(x.data_ptr<float>()),
            reinterpret_cast<float4*>(out.data_ptr<float>()),
            b.data_ptr<float>(), n4, C >> 2,
            (float)add_value, (float)scale);
    } else {
        long blocks = (n + block - 1) / block;
        if (blocks > 2147483647L) blocks = 2147483647L;
        fused_epilogue_nhwc_kernel<<<(unsigned int)blocks, block, 0, stream>>>(
            x.data_ptr<float>(), out.data_ptr<float>(), b.data_ptr<float>(),
            n, C, (float)add_value, (float)scale);
    }
    return out;
}
"""

fused_epilogue_cpp_source = (
    "torch::Tensor fused_epilogue_nhwc_cuda(torch::Tensor x, torch::Tensor bias, double add_value, double scale);"
)

fused_epilogue = load_inline(
    name="fused_bias_mish_add_hardtanh_scale_nhwc_v2",
    cpp_sources=fused_epilogue_cpp_source,
    cuda_sources=fused_epilogue_source,
    functions=["fused_epilogue_nhwc_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      * Runs the transposed convolution in channels_last so cudnn's fast NHWC
        dgrad GEMM is used without the expensive NHWC->NCHW output transpose.
      * Calls F.conv_transpose2d with bias=None and folds the bias into a
        single fused CUDA epilogue (bias + Mish + add + Hardtanh + scale),
        removing ATen's separate bias-add pass and three memory round-trips.
      * The conv_transpose module (incl. its bias parameter) is kept intact so
        the state_dict remains identical to the reference model.
    """
    def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale):
        super(ModelNew, self).__init__()
        self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride, padding, output_padding)
        self.add_value = add_value
        self.scale = scale
        self.fused_epilogue = fused_epilogue

    def forward(self, x):
        x = x.contiguous(memory_format=torch.channels_last)
        y = F.conv_transpose2d(
            x,
            self.conv_transpose.weight,
            None,  # bias folded into the fused epilogue
            self.conv_transpose.stride,
            self.conv_transpose.padding,
            self.conv_transpose.output_padding,
            self.conv_transpose.groups,
            self.conv_transpose.dilation,
        )
        y = y.contiguous(memory_format=torch.channels_last)
        bias = self.conv_transpose.bias
        if bias is None:
            bias = torch.zeros(y.size(1), device=y.device, dtype=y.dtype)
        y = self.fused_epilogue.fused_epilogue_nhwc_cuda(y, bias, self.add_value, self.scale)
        return y
```