message: Avoid ambiguous 'detail' namespace.
diff: @@ -641,20 +641,6 @@ Status LaunchDepthwiseConv2dGPUSmall(OpKernelContext* ctx,
  return Status::OK();
  }
  -namespace {
  -// Returns whether the context's GPU supports efficient fp16 math.
  -bool HasFastHalfMath(OpKernelContext* ctx) {
  - int major, minor;
  - ctx->op_device_context()
  - ->stream()
  - ->parent()
  - ->GetDeviceDescription()
  - .cuda_compute_capability(&major, &minor);
  - auto cuda_arch = major * 100 + minor * 10;
  - // GPUs before sm_53 don't support fp16 math, and sm_61's fp16 math is slow.
  - return cuda_arch >= 530 && cuda_arch != 610;
  -}
  -
  namespace detail {
  template <typename T>
  struct PseudoHalfType {
  @@ -666,9 +652,23 @@ struct PseudoHalfType<Eigen::half> {
  };
  } // namespace detail
  +namespace {
  // Maps to float if T is __half, and to T otherwise.
  template <typename T>
  using PseudoHalfType = typename detail::PseudoHalfType<T>::Type;
  +
  +// Returns whether the context's GPU supports efficient fp16 math.
  +bool HasFastHalfMath(OpKernelContext* ctx) {
  + int major, minor;
  + ctx->op_device_context()
  + ->stream()
  + ->parent()
  + ->GetDeviceDescription()
  + .cuda_compute_capability(&major, &minor);
  + auto cuda_arch = major * 100 + minor * 10;
  + // GPUs before sm_53 don't support fp16 math, and sm_61's fp16 math is slow.
  + return cuda_arch >= 530 && cuda_arch != 610;
  +}
  } // namespace
  template <typename T, DepthwiseConv2dDirection kDirection,
  