From 6e4ca3fbb4b8b4057ff9a636ffd2f1989126d097 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Tue, 28 Nov 2017 19:45:47 +0000
Subject: [PATCH] Change memset kernel to use memcpy instead of placement new.
 Simplify indexers.

---
 include/hip/hcc_detail/hip_runtime.h | 15 +++------------
 src/hip_memory.cpp                   |  5 ++++-
 2 files changed, 7 insertions(+), 13 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index c2302d4dc1..924e774af0 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -389,18 +389,9 @@ template<typename std::common_type<
 class Coordinates {
     using R = decltype(f(0));
 
-    struct X {
-        __device__ operator R() const { return f(0); }
-        __host__ operator R() const { return R{}; }
-    };
-    struct Y {
-        __device__ operator R() const { return f(1); }
-        __host__ operator R() const { return R{}; }
-    };
-    struct Z {
-        __device__ operator R() const { return f(2); }
-        __host__ operator R() const { return R{}; }
-    };
+    struct X { __device__ operator R() const { return f(0); } };
+    struct Y { __device__ operator R() const { return f(1); } };
+    struct Z { __device__ operator R() const { return f(2); } };
 public:
     static constexpr X x{};
     static constexpr Y y{};
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 502afa995b..c88a1dabc1 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1311,7 +1311,10 @@ namespace
 
         size_t idx = blockIdx.x * block_dim + threadIdx.x;
         while (idx < n) {
-            new (&f[idx]) T{value};
+            __builtin_memcpy(
+                reinterpret_cast<void*>(&f[idx]),
+                reinterpret_cast<const void*>(&value),
+                sizeof(T));
             idx += grid_dim;
         }
     }