ColumnReduceKernel: min() type casting error and improvement
I think following changes can solve the TODO:

[line 351-364](https://github.com/tensorflow/tensorflow/blob/master/tensorflow/core/kernels/reduction_gpu_kernels.cu.h#L351):
```
#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM
  __shared__ storage_type<value_type>
      partial_sums[TF_RED_WARPSIZE * (TF_RED_WARPSIZE + 1)];
#endif
```

[line 392](https://github.com/tensorflow/tensorflow/blob/master/tensorflow/core/kernels/reduction_gpu_kernels.cu.h#L392):
`min(blockDim.y, num_rows - blockIdx.y * blockDim.y); // MSVC type casting fix`

These changes can be compiled successfully in Windows CUDA environment. Can someone confirm if that's the meaning of the TODO? Hi @johnnkp 

Proposed PR has been merged now. Could you please check and close this issue.

Thank you!! The mentioned PR does not include my proposed changes. I am going to create a PR for this issue. Are you satisfied with the resolution of your issue?
<a href="https://docs.google.com/forms/d/e/1FAIpQLSfaP12TRhd9xSxjXZjcZFNXPGk4kc1-qMdv3gc6bEP90vY1ew/viewform?entry.85265664=Yes&entry.2137816233=https://github.com/tensorflow/tensorflow/issues/61357">Yes</a>
<a href="https://docs.google.com/forms/d/e/1FAIpQLSfaP12TRhd9xSxjXZjcZFNXPGk4kc1-qMdv3gc6bEP90vY1ew/viewform?entry.85265664=No&entry.2137816233=https://github.com/tensorflow/tensorflow/issues/61357">No</a>
 Hi @johnnkp 
```
min(blockDim.y, num_rows - blockIdx.y * blockDim.y); // MSVC type casting fix
```
this caused a lot of build errors due to ```call to 'min' is ambiguous```
https://github.com/tensorflow/tensorflow/pull/61638

### Issue type

Performance

### Have you reproduced the bug with TensorFlow Nightly?

No

### Source

source

### TensorFlow version

2.13.0

### Custom code

Yes

### OS platform and distribution

Windows 10 22H2

### Mobile device

_No response_

### Python version

Anaconda 2023.07-1

### Bazel version

6.2.1

### GCC/compiler version

Visual Studio 2022 (build tools 14.36) + msys2-x86_64-20230718

### CUDA/cuDNN version

CUDA 11.8 + CUDNN 8.6.0 + TensorRT 8.5.3

### GPU model and memory

GTX 750 Ti 2GB

### Current behavior?

There are two type casting errors in reduction_gpu_kernels.cu.h under MSVC. One of them is fixed in https://github.com/tensorflow/tensorflow/pull/61339. Another is related to a TODO.

in [ColumnReduceKernel()](https://github.com/tensorflow/tensorflow/blob/master/tensorflow/core/kernels/reduction_gpu_kernels.cu.h#L341), the TODO said the followings:

> 1D array necessary due to bug in CUDA 9 compiler.
> TODO(nluehr) revert to 2D array when compiler is ready.
> This is to mimic the following, but without constructors:
> __shared__ storage_type<value_type> partial_sums[TF_RED_WARPSIZE *
> (TF_RED_WARPSIZE + 1)];

Since latest version required CUDA 11, it's time to address the TODO and apply bug fix together.

### Standalone code to reproduce the issue

```shell
1. download https://github.com/tensorflow/tensorflow/archive/refs/tags/v2.13.0.zip and extract
2. comment out Windows CUDA build rejection code in configure.py
3. run `python configure.py` to configure Windows CUDA build
4. run `bazel build --config=opt --define=no_tensorflow_py_deps=true //tensorflow/tools/pip_package:build_pip_package`
```


### Relevant log output

```shell
external/com_google_absl\absl/status/status.h(796): warning #2810-D: ignoring return value type with "nodiscard" attribute

.\tensorflow/tsl/platform/file_system.h(571): warning #611-D: overloaded virtual function "tsl::FileSystem::FilesExist" is only partially overridden in class "tsl::WrappedFileSyste
m"

.\tensorflow/tsl/platform/file_system.h(571): warning #611-D: overloaded virtual function "tsl::FileSystem::CreateDir" is only partially overridden in class "tsl::WrappedFileSystem
"

.\tensorflow/tsl/platform/env.h(500): warning #611-D: overloaded virtual function "tsl::Env::RegisterFileSystem" is only partially overridden in class "tsl::EnvWrapper"

.\tensorflow/tsl/platform/float8.h(936): warning #177-D: variable "aligned_highest" was declared but never referenced
          detected during:
            instantiation of "To tsl::float8_internal::ConvertImpl<From, To, kSaturate, kTruncate, std::enable_if_t<<expression>, void>>::run(const From &) [with From=tsl::float8_i
nternal::float8_e4m3b11, To=tsl::float8_internal::float8_e4m3fn, kSaturate=false, kTruncate=false]"
(1018): here
            instantiation of "Derived tsl::float8_internal::float8_base<Derived>::ConvertFrom(const From &) [with Derived=tsl::float8_internal::float8_e4m3fn, kSaturate=false, kTru
ncate=false, From=tsl::float8_internal::float8_e4m3b11]"
(277): here

.\tensorflow/tsl/platform/float8.h(936): warning #177-D: variable "aligned_highest" was declared but never referenced
          detected during:
            instantiation of "To tsl::float8_internal::ConvertImpl<From, To, kSaturate, kTruncate, std::enable_if_t<<expression>, void>>::run(const From &) [with From=tsl::float8_i
nternal::float8_e4m3b11, To=float, kSaturate=false, kTruncate=false]"
(1024): here
            instantiation of "To tsl::float8_internal::float8_base<Derived>::ConvertTo<To,kSaturate,kTruncate>(const Derived &) [with Derived=tsl::float8_internal::float8_e4m3b11,
To=float, kSaturate=false, kTruncate=false]"
(75): here
            instantiation of "tsl::float8_internal::float8_base<Derived>::operator float() const [with Derived=tsl::float8_internal::float8_e4m3b11]"
(116): here
            instantiation of "Derived tsl::float8_internal::float8_base<Derived>::operator-(const Derived &) const [with Derived=tsl::float8_internal::float8_e4m3b11]"
(302): here

.\tensorflow/core/kernels/reduction_gpu_kernels.cu.h(392): error: no instance of overloaded function "tensorflow::min" matches the argument list
            argument types are: (int, unsigned int)
          detected during:
            instantiation of "void tensorflow::functor::ColumnReduceKernel(T, OUT_T, int, int, Op, std::iterator_traits<T>::value_type) [with T=const float *, OUT_T=float *, Op=cub
::Max]"
(828): here
            instantiation of "void tensorflow::functor::LaunchColumnReduction_LTE4096Cols(tensorflow::OpKernelContext *, OUT_T, IN_T, int, int, Op, T, const gpuStream_t &) [with T=
float, Op=cub::Max, OUT_T=float *, IN_T=const float *]"
(862): here
            instantiation of "void tensorflow::functor::LaunchColumnReduction(tensorflow::OpKernelContext *, OUT_T, IN_T, int, int, Op, T, const gpuStream_t &) [with T=float, Op=cu
b::Max, OUT_T=float *, IN_T=const float *]"
(1088): here
            instantiation of "void tensorflow::functor::ReduceImpl<T,Op,OUT_T,IN_T,ReductionAxes>(tensorflow::OpKernelContext *, OUT_T, IN_T, int, int, int, int, int, const Reducti
onAxes &, Op) [with T=float, Op=cub::Max, OUT_T=float *, IN_T=const float *, ReductionAxes=const Eigen::array<Eigen::DenseIndex, 1ULL> &]"
E:\_bazel_tensorflow\4zvk5ci6\execroot\org_tensorflow\tensorflow\core\kernels\multinomial_op_gpu.cu.cc(111): here
            instantiation of "void tensorflow::functor::MultinomialFunctor<tensorflow::functor::GPUDevice, T, OutputType>::operator()(tensorflow::OpKernelContext *, const tensorflo
w::functor::GPUDevice &, tensorflow::TTypes<T, 1, Eigen::DenseIndex>::ConstMatrix, tensorflow::TTypes<float, 1, Eigen::DenseIndex>::Flat, tensorflow::TTypes<float, 1, Eigen::DenseI
ndex>::Flat, tensorflow::TTypes<float, 1, Eigen::DenseIndex>::Flat, int, int, int, const tsl::random::PhiloxRandom &, tensorflow::TTypes<OutputType, 1, Eigen::DenseIndex>::Matrix)
[with T=Eigen::half, OutputType=tsl::int32]"
E:\_bazel_tensorflow\4zvk5ci6\execroot\org_tensorflow\tensorflow\core\kernels\multinomial_op_gpu.cu.cc(126): here

1 error detected in the compilation of "tensorflow/core/kernels/multinomial_op_gpu.cu.cc".
nvcc warning : The 'compute_35', 'compute_37', 'sm_35', and 'sm_37' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppres
s warning).
Target //tensorflow/tools/pip_package:build_pip_package failed to build
INFO: Elapsed time: 1996.828s, Critical Path: 480.15s
INFO: 441 processes: 7 internal, 434 local.
FAILED: Build did NOT complete successfully
```
