```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 (improvements over the 5.14x turn-3 solution):
#   * Same algorithmic rewrite: stride-1 ConvTranspose3d == forward Conv3d
#     with flipped/channel-swapped weights and padding (K-1-p) (cached).
#   * Vectorized custom pack/unpack kernels (half2 / float4 global accesses)
#     for the fp32-NCHW <-> fp16-NHWC layout+cast transforms around the fp16
#     tensor-core convolution.
#   * New batch-split software pipeline: the fp16 conv is split into two half
#     batches on the main stream while the pack of the second half and the
#     unpack of the first half run concurrently on a side stream, hiding most
#     of the transpose overhead behind the compute-bound convolution.
#   * A third candidate path (fp16 NCHW conv with cheap elementwise casts)
#     plus an exhaustive cuDNN engine search (benchmark_limit = 0, scoped to
#     the one-time autotune).  A CUDA-event autotune picks the fastest of
#     {pipelined fp16 NHWC, serial fp16 NHWC, fp16 NCHW, fp32 TF32 NCHW} on
#     the actual hardware, so the solution can never regress below the
#     proven fp32 path.
# ---------------------------------------------------------------------------

_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);
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)
// mode 2: fp16 NCHW contiguous output
// -------------------------------------------------------------------------
__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_nchw_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 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] = __float2half(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 * (TILE * TILE / 2) / 2;  // 256 threads, 512 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;
  }
  if (mode == 2) {
    auto out = at::empty({C_out, C_in_g, KD, KH, KW},
                         w_c.options().dtype(at::kHalf));
    flip_weight_fp16_nchw_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) {
  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;
  auto 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_v4",
        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, 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 paths, chosen by a one-time on-device autotune:
      * batch-split pipelined fp16 NHWC tensor-core conv whose pack/unpack
        layout transforms overlap the compute on a second CUDA stream,
      * serial fp16 NHWC conv with the custom fused transform kernels,
      * fp16 NCHW conv with plain elementwise casts,
      * fp32 (TF32) NCHW conv.
    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._cache_key = None
        self._w2 = None       # flipped fp32 NCHW weight (cached)
        self._w_h = None      # flipped fp16 NHWC weight (cached)
        self._w_h2 = None     # flipped fp16 NCHW weight (cached)
        self._bias_h = None   # fp16 bias (cached)
        self._mode = None     # None -> autotune on first fast call
        self._stream = None
        self._ev_p2 = None
        self._ev_c1 = None
        self._ev_u1 = None

    def _ensure_stream(self):
        if self._stream is None:
            self._stream = torch.cuda.Stream()
            self._ev_p2 = torch.cuda.Event()
            self._ev_c1 = torch.cuda.Event()
            self._ev_u1 = torch.cuda.Event()

    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)
            except Exception:
                self._w_h = None
            try:
                self._w_h2 = _ext.flip_weight(w, g, 2)
            except Exception:
                self._w_h2 = None
            try:
                self._bias_h = (bias.to(torch.float16)
                                if bias is not None else None)
            except Exception:
                self._bias_h = None
            self._cache_key = key
            self._mode = None

    def _conv_args(self):
        ct = self.conv_transpose3d
        return ((1, 1, 1), self._conv_pad, (1, 1, 1), ct.groups)

    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)
        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 _run_fp16_nchw(self, x):
        y = F.conv3d(x.half(), self._w_h2, self._bias_h, *self._conv_args())
        y = y.float()
        if not y.is_contiguous():
            y = y.contiguous()
        return y

    def _run_fp16_pipe(self, x):
        n = x.size(0)
        if n < 2:
            return self._run_fp16(x)
        self._ensure_stream()
        n1 = (n + 1) // 2
        x1 = x[:n1]
        x2 = x[n1:]
        cur = torch.cuda.current_stream()
        s2 = self._stream
        s2.wait_stream(cur)                      # x is produced on cur
        x1_h = _ext.pack_half(x1)                # main stream
        with torch.cuda.stream(s2):
            x2_h = _ext.pack_half(x2)            # overlaps conv1
            self._ev_p2.record(s2)
        y1 = F.conv3d(x1_h, self._w_h, self._bias_h, *self._conv_args())
        self._ev_c1.record(cur)
        if not y1.is_contiguous(torch.channels_last_3d):
            raise RuntimeError("unexpected conv output layout")
        cur.wait_event(self._ev_p2)
        y2 = F.conv3d(x2_h, self._w_h, self._bias_h, *self._conv_args())
        if not y2.is_contiguous(torch.channels_last_3d):
            raise RuntimeError("unexpected conv output layout")
        out = torch.empty((n,) + tuple(y1.shape[1:]), dtype=torch.float32,
                          device=x.device)
        with torch.cuda.stream(s2):
            s2.wait_event(self._ev_c1)
            _ext.unpack_half(y1, out[:n1])       # overlaps conv2
            self._ev_u1.record(s2)
        _ext.unpack_half(y2, out[n1:])
        cur.wait_event(self._ev_u1)
        # Tell the caching allocator about cross-stream tensor uses.
        x.record_stream(s2)
        x2_h.record_stream(cur)
        y1.record_stream(s2)
        out.record_stream(s2)
        return out

    def _autotune(self, x):
        candidates = []
        if self._w_h is not None and x.size(0) >= 2:
            candidates.append(("fp16_pipe", self._run_fp16_pipe))
        if self._w_h is not None:
            candidates.append(("fp16_nhwc", self._run_fp16))
        if self._w_h2 is not None:
            candidates.append(("fp16_nchw", self._run_fp16_nchw))
        candidates.append(("fp32", self._run_fp32))
        old_limit = torch.backends.cudnn.benchmark_limit
        try:
            torch.backends.cudnn.benchmark_limit = 0  # exhaustive engine search
            valid = []
            for name, fn in candidates:
                try:
                    fn(x)
                    fn(x)  # warm up cudnn benchmarking for this config
                    valid.append((name, fn))
                except Exception:
                    pass
            if not valid:
                self._mode = "native"
                return
            if len(valid) == 1:
                self._mode = valid[0][0]
                return
            torch.cuda.synchronize()
            best, best_t = None, 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 if best is not None else "native"
        except Exception:
            self._mode = "fp32"
        finally:
            torch.backends.cudnn.benchmark_limit = old_limit

    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 == "fp16_pipe":
                    try:
                        return self._run_fp16_pipe(x)
                    except Exception:
                        m = self._mode = ("fp16_nhwc"
                                          if self._w_h is not None else "fp32")
                if m == "fp16_nhwc" and self._w_h is not None:
                    try:
                        return self._run_fp16(x)
                    except Exception:
                        m = self._mode = "fp32"
                if m == "fp16_nchw" and self._w_h2 is not None:
                    try:
                        return self._run_fp16_nchw(x)
                    except Exception:
                        m = self._mode = "fp32"
                if m == "fp32" and self._w2 is not None:
                    return self._run_fp32(x)
            except Exception:
                self._mode = "native"
        return ct(x)
```