[release/2.9] Fix int4mm device memcpy error on Windows (pytorch#175410) (#3166)

On Windows with HIP/ROCm, std::memcpy is a __host__ function and cannot
be called from __device__ code. Use raw memcpy (which the HIP compiler
provides as a device builtin) when building on Windows.

This will allow builds for of pytorch for gfx942 on Windows. gfx950 is
yet to be tested but it should likely build as well.

Pull Request resolved: pytorch#175410
Approved by: https://github.com/jeffdaily

Co-authored-by: Aaryaman Vasishta <aaryaman.vasishta@amd.com>

Author: jeffdaily

aten/src/ATen/native/cuda/int4mm.cu

@@ -581,7 +581,14 @@ struct BLayout_TC_int4 {
           // type pun, the __nv_bfloat162 value in bf16x2x4 is a struct and
           // can't be used as a 32-bit asm register argument for `mma`
           static_assert(sizeof(bf16x2x4) == sizeof(out[0][0]), "");
+          // On Windows with ROCm, std::memcpy resolves to a __host__-only
+          // function and cannot be called from __device__ code. Use the raw
+          // memcpy which the HIP compiler provides as a __device__ builtin.
+#if defined(_WIN32) && defined(USE_ROCM)
+          memcpy(&out[i][j], &v, sizeof(bf16x2x4_u32));
+#else
           std::memcpy(&out[i][j], &v, sizeof(bf16x2x4_u32));
+#endif
         }
       }
     }