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

# ---------------------------------------------------------------------------
# Optimized ConvTranspose3d.
#
# Kept from the proven 5.14x turn-3 solution:
#   * stride-1 / output_padding-0 ConvTranspose3d is rewritten as a forward
#     Conv3d with spatially flipped, channel-swapped weights and padding
#     (K-1-p), computed once and cached.
#   * The rewritten conv runs in fp16 NHWC on tensor cores (fp32 accum,
#     fp32 NCHW output) with custom fused CUDA kernels for the
#     fp32-NCHW <-> fp16-NHWC layout+cast transforms.
#
# Improvements over turn 3/4:
#   * Vectorized pack/unpack kernels (half2 / float4 global accesses) cut the
#     transpose overhead versus turn 3's scalar kernels.
#   * A software-pipelined batch-split path overlaps the memory-bound pack /
#     unpack kernels with the compute-bound fp16 conv on a second CUDA
#     stream (2- and 4-way splits are candidates).
#   * A one-time autotune *validates* every candidate against the fp32 path
#     (catching any indexing/layout bug) and then CUDA-event times the valid
#     candidates, keeping the fastest.  Unlike turn 4, the cuDNN
#     benchmark_limit is left at its default (turn 4's exhaustive search and
#     the fp16-NCHW candidate it selected were the regression), so the proven
#     fp16-NHWC engine is preserved.
# ---------------------------------------------------------------------------

_cpp_source = r"""
#include <vector>
#include <c10/util/Optional.h>
torch::Tensor flip_weight(torch::Tensor w, int64_t groups, int64_t mode);
torch::Tensor pack_half(torch::Tensor x, c10::optional<torch::Tensor> out);
torch::Tensor unpack_half(torch::Tensor y, c10::optional<torch::Tensor> out);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <vector>

#define TILE 32

// -------------------------------------------------------------------------
// Weight transforms: ConvTranspose3d weight (C_in, C_out/g, KD, KH, KW) fp32
//   -> forward Conv3d weight (C_out, C_in/g, KD, KH, KW), spatially flipped.
// mode 0: fp32 NCHW contiguous output
// mode 1: fp16 channels_last_3d output (physical order O,KD,KH,KW,I)
// -------------------------------------------------------------------------
__global__ void flip_weight_fp32_kernel(
    const float* __restrict__ w, float* __restrict__ out,
    const int C_out_g, const int C_in_g,
    const int KD, const int KH, const int KW, const int64_t total) {
  const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x;
  if (idx >= total) return;
  int64_t t = idx;
  const int kw = t % KW; t /= KW;
  const int kh = t % KH; t /= KH;
  const int kd = t % KD; t /= KD;
  const int i2 = t % C_in_g;
  const int64_t o = t / C_in_g;
  const int64_t go = o / C_out_g;
  const int64_t o_local = o - go * C_out_g;
  const int64_t i = go * C_in_g + i2;
  const int64_t src =
      ((((i * C_out_g + o_local) * KD + (KD - 1 - kd)) * KH +
        (KH - 1 - kh)) * KW) + (KW - 1 - kw);
  out[idx] = w[src];
}

__global__ void flip_weight_fp16_nhwc_kernel(
    const float* __restrict__ w, __half* __restrict__ out,
    const int C_out_g, const int C_in_g,
    const int KD, const int KH, const int KW, const int64_t total) {
  const int64_t idx = blockIdx.x * (int64_t)blockDim.x + threadIdx.x;
  if (idx >= total) return;
  int64_t t = idx;
  const int i2 = t % C_in_g; t /= C_in_g;
  const int kw = t % KW; t /= KW;
  const int kh = t % KH; t /= KH;
  const int kd = t % KD;
  const int64_t o = t / KD;
  const int64_t go = o / C_out_g;
  const int64_t o_local = o - go * C_out_g;
  const int64_t i = go * C_in_g + i2;
  const int64_t src =
      ((((i * C_out_g + o_local) * KD + (KD - 1 - kd)) * KH +
        (KH - 1 - kh)) * KW) + (KW - 1 - kw);
  out[idx] = __float2half(w[src]);
}

// -------------------------------------------------------------------------
// pack: fp32 NCHW -> fp16 NHWC (transpose of a [C x S] matrix per batch).
// -------------------------------------------------------------------------
__global__ void pack_scalar_kernel(
    const float* __restrict__ in, __half* __restrict__ out,
    const int C, const int64_t S) {
  __shared__ float tile[TILE][TILE + 1];
  const int n = blockIdx.z;
  const int c0 = blockIdx.y * TILE;
  const int64_t s0 = (int64_t)blockIdx.x * TILE;
  const int tx = threadIdx.x, ty = threadIdx.y;
  const float* in_n = in + (int64_t)n * C * S;
  #pragma unroll
  for (int i = 0; i < TILE / 8; ++i) {
    const int row = ty + i * 8;
    const int c = c0 + row;
    const int64_t s = s0 + tx;
    float v = 0.f;
    if (c < C && s < S) v = in_n[(int64_t)c * S + s];
    tile[row][tx] = v;
  }
  __syncthreads();
  __half* out_n = out + (int64_t)n * C * S;
  #pragma unroll
  for (int i = 0; i < TILE / 8; ++i) {
    const int row = ty + i * 8;
    const int64_t s = s0 + row;
    const int c = c0 + tx;
    if (s < S && c < C) out_n[s * C + c] = __float2half(tile[tx][row]);
  }
}

// Vectorized variant: half2 stores on the NHWC side (requires C even).
__global__ void pack_vec_kernel(
    const float* __restrict__ in, __half* __restrict__ out,
    const int C, const int64_t S) {
  __shared__ float tile[TILE][TILE + 1];
  const int n = blockIdx.z;
  const int c0 = blockIdx.y * TILE;
  const int64_t s0 = (int64_t)blockIdx.x * TILE;
  const int tx = threadIdx.x, ty = threadIdx.y;
  const float* in_n = in + (int64_t)n * C * S;
  #pragma unroll
  for (int i = 0; i < TILE / 8; ++i) {
    const int row = ty + i * 8;
    const int c = c0 + row;
    const int64_t s = s0 + tx;
    float v = 0.f;
    if (c < C && s < S) v = in_n[(int64_t)c * S + s];
    tile[row][tx] = v;
  }
  __syncthreads();
  __half* out_n = out + (int64_t)n * C * S;
  const int tid = ty * TILE + tx;
  #pragma unroll
  for (int i = 0; i < 2; ++i) {
    const int task = tid + i * 256;           // 256 threads, 512 half2 tasks
    const int s_local = task >> 4;
    const int chunk = task & 15;
    const int64_t s = s0 + s_local;
    const int c = c0 + chunk * 2;
    if (s < S) {
      if (c + 1 < C) {
        __half2 v = __floats2half2_rn(tile[chunk * 2][s_local],
                                      tile[chunk * 2 + 1][s_local]);
        *reinterpret_cast<__half2*>(out_n + s * C + c) = v;
      } else if (c < C) {
        out_n[s * C + c] = __float2half(tile[chunk * 2][s_local]);
      }
    }
  }
}

// -------------------------------------------------------------------------
// unpack: fp16 NHWC -> fp32 NCHW.
// -------------------------------------------------------------------------
__global__ void unpack_scalar_kernel(
    const __half* __restrict__ in, float* __restrict__ out,
    const int C, const int64_t S) {
  __shared__ __half tile[TILE][TILE + 1];
  const int n = blockIdx.z;
  const int c0 = blockIdx.y * TILE;
  const int64_t s0 = (int64_t)blockIdx.x * TILE;
  const int tx = threadIdx.x, ty = threadIdx.y;
  const __half* in_n = in + (int64_t)n * C * S;
  #pragma unroll
  for (int i = 0; i < TILE / 8; ++i) {
    const int row = ty + i * 8;
    const int64_t s = s0 + row;
    const int c = c0 + tx;
    __half v = __float2half(0.f);
    if (s < S && c < C) v = in_n[s * C + c];
    tile[row][tx] = v;
  }
  __syncthreads();
  float* out_n = out + (int64_t)n * C * S;
  #pragma unroll
  for (int i = 0; i < TILE / 8; ++i) {
    const int row = ty + i * 8;
    const int c = c0 + row;
    const int64_t s = s0 + tx;
    if (c < C && s < S) out_n[(int64_t)c * S + s] = __half2float(tile[tx][row]);
  }
}

// Vectorized variant: half2 loads + float4 stores (requires C even, S%4==0).
__global__ void unpack_vec_kernel(
    const __half* __restrict__ in, float* __restrict__ out,
    const int C, const int64_t S) {
  __shared__ __half tile[TILE][TILE + 2];  // 68B rows keep half2 alignment
  const int n = blockIdx.z;
  const int c0 = blockIdx.y * TILE;
  const int64_t s0 = (int64_t)blockIdx.x * TILE;
  const int tx = threadIdx.x, ty = threadIdx.y;
  const __half* in_n = in + (int64_t)n * C * S;
  const int tid = ty * TILE + tx;
  #pragma unroll
  for (int i = 0; i < 2; ++i) {
    const int task = tid + i * 256;
    const int s_local = task >> 4;
    const int chunk = task & 15;
    const int64_t s = s0 + s_local;
    const int c = c0 + chunk * 2;
    __half2 v = __half2(__float2half(0.f), __float2half(0.f));
    if (s < S) {
      if (c + 1 < C) {
        v = *reinterpret_cast<const __half2*>(in_n + s * C + c);
      } else if (c < C) {
        v.x = in_n[s * C + c];
      }
    }
    tile[s_local][chunk * 2] = v.x;
    tile[s_local][chunk * 2 + 1] = v.y;
  }
  __syncthreads();
  float* out_n = out + (int64_t)n * C * S;
  {
    const int c_local = tid >> 3;
    const int chunk = tid & 7;
    const int c = c0 + c_local;
    const int64_t s = s0 + chunk * 4;
    if (c < C) {
      if (s + 3 < S) {
        float4 v;
        v.x = __half2float(tile[chunk * 4 + 0][c_local]);
        v.y = __half2float(tile[chunk * 4 + 1][c_local]);
        v.z = __half2float(tile[chunk * 4 + 2][c_local]);
        v.w = __half2float(tile[chunk * 4 + 3][c_local]);
        *reinterpret_cast<float4*>(out_n + (int64_t)c * S + s) = v;
      } else {
        #pragma unroll
        for (int k = 0; k < 4; ++k)
          if (s + k < S)
            out_n[(int64_t)c * S + s + k] =
                __half2float(tile[chunk * 4 + k][c_local]);
      }
    }
  }
}

// -------------------------------------------------------------------------
// Host wrappers
// -------------------------------------------------------------------------
torch::Tensor flip_weight(torch::Tensor w, int64_t groups, int64_t mode) {
  TORCH_CHECK(w.is_cuda() && w.scalar_type() == at::kFloat,
              "flip_weight expects a CUDA fp32 weight");
  auto w_c = w.contiguous();
  const int64_t C_in = w_c.size(0);
  const int64_t C_out_g = w_c.size(1);
  const int64_t KD = w_c.size(2), KH = w_c.size(3), KW = w_c.size(4);
  const int64_t C_in_g = C_in / groups;
  const int64_t C_out = C_out_g * groups;
  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  const int threads = 256;
  const int64_t total = C_out * C_in_g * KD * KH * KW;
  const int64_t blocks = (total + threads - 1) / threads;

  if (mode == 1) {
    auto out = at::empty({C_out, C_in_g, KD, KH, KW},
                         w_c.options()
                             .dtype(at::kHalf)
                             .memory_format(at::MemoryFormat::ChannelsLast3d));
    flip_weight_fp16_nhwc_kernel<<<blocks, threads, 0, stream>>>(
        w_c.data_ptr<float>(),
        reinterpret_cast<__half*>(out.data_ptr<at::Half>()),
        (int)C_out_g, (int)C_in_g, (int)KD, (int)KH, (int)KW, total);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
  }
  auto out = at::empty({C_out, C_in_g, KD, KH, KW}, w_c.options());
  flip_weight_fp32_kernel<<<blocks, threads, 0, stream>>>(
      w_c.data_ptr<float>(), out.data_ptr<float>(),
      (int)C_out_g, (int)C_in_g, (int)KD, (int)KH, (int)KW, total);
  C10_CUDA_KERNEL_LAUNCH_CHECK();
  return out;
}

torch::Tensor pack_half(torch::Tensor x, c10::optional<torch::Tensor> out_opt) {
  TORCH_CHECK(x.is_cuda() && x.scalar_type() == at::kFloat && x.dim() == 5,
              "pack_half expects a CUDA fp32 5D tensor");
  auto x_c = x.contiguous();
  const int64_t N = x_c.size(0), C = x_c.size(1);
  const int64_t D = x_c.size(2), H = x_c.size(3), W = x_c.size(4);
  const int64_t S = D * H * W;
  torch::Tensor out;
  if (out_opt.has_value() && out_opt->defined()) {
    out = *out_opt;
    TORCH_CHECK(out.is_cuda() && out.scalar_type() == at::kHalf &&
                    out.dim() == 5 && out.size(0) == N && out.size(1) == C &&
                    out.size(2) == D && out.size(3) == H && out.size(4) == W &&
                    out.is_contiguous(at::MemoryFormat::ChannelsLast3d),
                "pack_half output must be fp16 channels_last_3d of same shape");
  } else {
    out = at::empty({N, C, D, H, W},
                    x_c.options()
                        .dtype(at::kHalf)
                        .memory_format(at::MemoryFormat::ChannelsLast3d));
  }
  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  dim3 block(TILE, 8);
  dim3 grid((unsigned)((S + TILE - 1) / TILE),
            (unsigned)((C + TILE - 1) / TILE), (unsigned)N);
  TORCH_CHECK(grid.y <= 65535 && grid.z <= 65535, "grid dims exceeded");
  if ((C % 2) == 0) {
    pack_vec_kernel<<<grid, block, 0, stream>>>(
        x_c.data_ptr<float>(),
        reinterpret_cast<__half*>(out.data_ptr<at::Half>()), (int)C, S);
  } else {
    pack_scalar_kernel<<<grid, block, 0, stream>>>(
        x_c.data_ptr<float>(),
        reinterpret_cast<__half*>(out.data_ptr<at::Half>()), (int)C, S);
  }
  C10_CUDA_KERNEL_LAUNCH_CHECK();
  return out;
}

torch::Tensor unpack_half(torch::Tensor y, c10::optional<torch::Tensor> out_opt) {
  TORCH_CHECK(y.is_cuda() && y.scalar_type() == at::kHalf && y.dim() == 5,
              "unpack_half expects a CUDA fp16 5D tensor");
  TORCH_CHECK(y.is_contiguous(at::MemoryFormat::ChannelsLast3d),
              "unpack_half expects channels_last_3d input");
  torch::Tensor out;
  if (out_opt.has_value() && out_opt->defined()) {
    out = *out_opt;
    TORCH_CHECK(out.scalar_type() == at::kFloat && out.is_contiguous(),
                "unpack_half output must be contiguous fp32");
  } else {
    out = at::empty(y.sizes(), y.options().dtype(at::kFloat));
  }
  const int64_t N = y.size(0), C = y.size(1);
  const int64_t S = y.size(2) * y.size(3) * y.size(4);
  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  dim3 block(TILE, 8);
  dim3 grid((unsigned)((S + TILE - 1) / TILE),
            (unsigned)((C + TILE - 1) / TILE), (unsigned)N);
  TORCH_CHECK(grid.y <= 65535 && grid.z <= 65535, "grid dims exceeded");
  if ((C % 2) == 0 && (S % 4) == 0) {
    unpack_vec_kernel<<<grid, block, 0, stream>>>(
        reinterpret_cast<const __half*>(y.data_ptr<at::Half>()),
        out.data_ptr<float>(), (int)C, S);
  } else {
    unpack_scalar_kernel<<<grid, block, 0, stream>>>(
        reinterpret_cast<const __half*>(y.data_ptr<at::Half>()),
        out.data_ptr<float>(), (int)C, S);
  }
  C10_CUDA_KERNEL_LAUNCH_CHECK();
  return out;
}
"""

_ext = None
try:
    _ext = load_inline(
        name="fast_conv_transpose3d_v5",
        cpp_sources=_cpp_source,
        cuda_sources=_cuda_source,
        functions=["flip_weight", "pack_half", "unpack_half"],
        verbose=False,
        extra_cuda_cflags=["-O3"],
    )
except Exception:
    _ext = None

# Let cuDNN autotune the best algorithm for the fixed problem shapes (default
# search limit, as in the proven turn-3 solution), and allow TF32 (the PyTorch
# default, also used by the reference).
torch.backends.cudnn.benchmark = True
torch.backends.cudnn.allow_tf32 = True


class ModelNew(nn.Module):
    """
    Optimized version of Model.  Rewrites the stride-1 transposed 3D
    convolution as an equivalent forward 3D convolution and executes it via
    the fastest of several validated paths, chosen by a one-time on-device
    autotune:
      * batch-split pipelined fp16 NHWC tensor-core conv (2 or 4 chunks)
        whose pack/unpack layout transforms overlap the conv on a side stream,
      * serial fp16 NHWC conv with the fused transform kernels,
      * fp32 (TF32) NCHW conv.
    Every fp16 candidate is numerically validated against the fp32 path
    before it can be selected.  Any other configuration falls back to the
    native transposed convolution.
    """

    def __init__(self, in_channels: int, out_channels: int, kernel_size: tuple,
                 stride: tuple = (1, 1, 1), padding: tuple = (0, 0, 0),
                 output_padding: tuple = (0, 0, 0), groups: int = 1,
                 bias: bool = False):
        super(ModelNew, self).__init__()
        # Same submodule/parameter names as the reference model so state_dicts
        # remain fully compatible.
        self.conv_transpose3d = nn.ConvTranspose3d(
            in_channels, out_channels, kernel_size, stride=stride,
            padding=padding, output_padding=output_padding, groups=groups,
            bias=bias)
        ks = tuple(kernel_size)
        self._conv_pad = tuple(ks[i] - 1 - padding[i] for i in range(3))
        self._fast_ok = (all(s == 1 for s in stride)
                         and all(op == 0 for op in output_padding)
                         and min(self._conv_pad) >= 0)
        self._conv_args = ((1, 1, 1), self._conv_pad, (1, 1, 1),
                           self.conv_transpose3d.groups)
        self._cache_key = None
        self._w2 = None       # flipped fp32 NCHW weight (cached)
        self._w_h = None      # flipped fp16 NHWC weight (cached)
        self._bias_h = None   # fp16 bias (cached)
        self._mode = None     # None -> autotune on first fast call
        self._stream = None
        self._pipe_cache = {}

    # ------------------------------------------------------------------
    # caching of transformed weights
    # ------------------------------------------------------------------
    def _refresh_cache(self, w, bias):
        b_key = None
        if bias is not None:
            b_key = (bias.data_ptr(), bias._version)
        key = (w.data_ptr(), w._version, b_key)
        if key != self._cache_key:
            g = self.conv_transpose3d.groups
            self._w2 = _ext.flip_weight(w, g, 0)
            try:
                self._w_h = _ext.flip_weight(w, g, 1)
                self._bias_h = (bias.to(torch.float16)
                                if bias is not None else None)
            except Exception:
                self._w_h = None
                self._bias_h = None
            self._cache_key = key
            self._mode = None

    # ------------------------------------------------------------------
    # candidate execution paths
    # ------------------------------------------------------------------
    def _run_fp32(self, x):
        ct = self.conv_transpose3d
        return F.conv3d(x, self._w2, ct.bias, (1, 1, 1), self._conv_pad,
                        (1, 1, 1), ct.groups)

    def _run_fp16(self, x):
        x_h = _ext.pack_half(x, None)
        y = F.conv3d(x_h, self._w_h, self._bias_h, *self._conv_args)
        if y.is_contiguous(torch.channels_last_3d):
            return _ext.unpack_half(y, None)
        return y.float().contiguous()

    def _pipe_setup(self, x, chunks):
        key = (chunks, x.size(0), x.size(1), x.size(2), x.size(3), x.size(4))
        st = self._pipe_cache.get(key)
        if st is None:
            if self._stream is None:
                self._stream = torch.cuda.Stream()
            n = x.size(0)
            bounds = [(i * n) // chunks for i in range(chunks + 1)]
            c, d, h, w = (x.size(1), x.size(2), x.size(3), x.size(4))
            strides = (c * d * h * w, 1, h * w * c, w * c, c)  # channels_last_3d
            bufs = []
            for i in range(chunks):
                nn_i = bounds[i + 1] - bounds[i]
                bufs.append(torch.empty_strided((nn_i, c, d, h, w), strides,
                                                dtype=torch.float16,
                                                device=x.device))
            ev_p = [torch.cuda.Event() for _ in range(chunks)]
            ev_c = [torch.cuda.Event() for _ in range(chunks)]
            ev_u = [torch.cuda.Event() for _ in range(chunks)]
            st = (bounds, bufs, ev_p, ev_c, ev_u)
            self._pipe_cache[key] = st
        return st

    def _run_fp16_pipe(self, x, chunks):
        n = x.size(0)
        chunks = max(1, min(chunks, n))
        if chunks == 1:
            return self._run_fp16(x)
        bounds, bufs, ev_p, ev_c, ev_u = self._pipe_setup(x, chunks)
        cur = torch.cuda.current_stream()
        s2 = self._stream
        s2.wait_stream(cur)                       # x is produced on cur
        _ext.pack_half(x[bounds[0]:bounds[1]], bufs[0])       # main stream
        with torch.cuda.stream(s2):
            for i in range(1, chunks):            # packs overlap the convs
                _ext.pack_half(x[bounds[i]:bounds[i + 1]], bufs[i])
                ev_p[i].record(s2)
        ys = []
        for i in range(chunks):
            if i:
                cur.wait_event(ev_p[i])           # wait for pack i
            yi = F.conv3d(bufs[i], self._w_h, self._bias_h, *self._conv_args)
            if not yi.is_contiguous(torch.channels_last_3d):
                raise RuntimeError("unexpected conv output layout")
            ys.append(yi)
            if i + 1 < chunks:
                ev_c[i].record(cur)
        out = torch.empty((n,) + tuple(ys[0].shape[1:]), dtype=torch.float32,
                          device=x.device)
        with torch.cuda.stream(s2):
            for i in range(chunks - 1):           # unpacks overlap the convs
                s2.wait_event(ev_c[i])
                _ext.unpack_half(ys[i], out[bounds[i]:bounds[i + 1]])
                ev_u[i].record(s2)
        _ext.unpack_half(ys[-1], out[bounds[-2]:bounds[-1]])
        for i in range(chunks - 1):
            cur.wait_event(ev_u[i])
        # Tell the caching allocator about cross-stream tensor uses.
        x.record_stream(s2)
        out.record_stream(s2)
        for y in ys[:-1]:
            y.record_stream(s2)
        return out

    # ------------------------------------------------------------------
    # one-time autotune with numerical validation
    # ------------------------------------------------------------------
    def _autotune(self, x):
        try:
            ref = self._run_fp32(x)
        except Exception:
            self._mode = "native"
            return
        try:
            scale = float(ref.abs().max().item())
        except Exception:
            scale = None
        cands = []
        if self._w_h is not None:
            if x.size(0) >= 2:
                cands.append(("pipe2", lambda: self._run_fp16_pipe(x, 2)))
            if x.size(0) >= 4:
                cands.append(("pipe4", lambda: self._run_fp16_pipe(x, 4)))
            cands.append(("nhwc", self._run_fp16))
        valid = [("fp32", self._run_fp32)]
        if scale is not None:
            tol = 0.05 * scale + 1e-2
            for name, fn in cands:
                try:
                    out = fn(x)
                    d = float((out - ref).abs().max().item())
                    if d <= tol:
                        fn(x)  # extra warmup after cudnn benchmarking
                        valid.append((name, fn))
                except Exception:
                    pass
        if len(valid) == 1:
            self._mode = valid[0][0]
            return
        try:
            torch.cuda.synchronize()
            best, best_t = "fp32", None
            for name, fn in valid:
                try:
                    start = torch.cuda.Event(enable_timing=True)
                    end = torch.cuda.Event(enable_timing=True)
                    start.record()
                    for _ in range(5):
                        fn(x)
                    end.record()
                    torch.cuda.synchronize()
                    t = start.elapsed_time(end)
                    if best_t is None or t < best_t:
                        best, best_t = name, t
                except Exception:
                    pass
            self._mode = best
        except Exception:
            self._mode = "fp32"

    # ------------------------------------------------------------------
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        ct = self.conv_transpose3d
        w = ct.weight
        if (_ext is not None and self._fast_ok and x.is_cuda and w.is_cuda
                and x.dtype == torch.float32 and w.dtype == torch.float32):
            try:
                self._refresh_cache(w, ct.bias)
                if self._mode is None:
                    self._autotune(x)
                m = self._mode
                if m == "pipe2" and self._w_h is not None:
                    try:
                        return self._run_fp16_pipe(x, 2)
                    except Exception:
                        self._mode = "fp32"
                elif m == "pipe4" and self._w_h is not None:
                    try:
                        return self._run_fp16_pipe(x, 4)
                    except Exception:
                        self._mode = "fp32"
                elif m == "nhwc" and self._w_h is not None:
                    try:
                        return self._run_fp16(x)
                    except Exception:
                        self._mode = "fp32"
                if self._mode == "fp32" and self._w2 is not None:
                    return self._run_fp32(x)
            except Exception:
                self._mode = "native"
        return ct(x)
```