From cc0fbc0e2ebbe3ce71347b3f00a586bf5d87da66 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Tue, 28 Aug 2018 19:13:27 +0000
Subject: [PATCH] Support placement new in hip-clang

---
 hipamd/include/hip/hcc_detail/hip_runtime.h |  1 +
 hipamd/tests/src/deviceLib/hipTestNew.cpp   | 71 +++++++++++++++++++++
 2 files changed, 72 insertions(+)
 create mode 100644 hipamd/tests/src/deviceLib/hipTestNew.cpp

diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index 41710e4165..7bd89e2815 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -471,6 +471,7 @@ hc_get_workitem_absolute_id(int dim)
 #include <__clang_cuda_complex_builtins.h>
 #include <cuda_wrappers/algorithm>
 #include <cuda_wrappers/complex>
+#include <cuda_wrappers/new>
 #undef __CUDA__
 #pragma pop_macro("__CUDA__")
 
diff --git a/hipamd/tests/src/deviceLib/hipTestNew.cpp b/hipamd/tests/src/deviceLib/hipTestNew.cpp
new file mode 100644
index 0000000000..60774ff21d
--- /dev/null
+++ b/hipamd/tests/src/deviceLib/hipTestNew.cpp
@@ -0,0 +1,71 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#include <iostream>
+#include "test_common.h"
+
+#define HIP_ASSERT(status) assert(status == hipSuccess)
+
+#define LEN 512
+#define SIZE 2048
+
+struct TestPlacementNew {
+  class A {
+  public:
+      __device__ A() {
+          a = threadIdx.x + blockIdx.x * blockDim.x;
+      }
+  private:
+      int a;
+  };
+     
+  static __global__ void kernel(int* Ad) {
+      int tid = threadIdx.x + blockIdx.x * blockDim.x;
+      new(Ad+tid) A();
+  }
+
+  void run() {
+    int *A, *Ad;
+    A = new int[LEN];
+    for (unsigned i = 0; i < LEN; i++) {
+        A[i] = 0;
+    }
+
+    HIP_ASSERT(hipMalloc((void**)&Ad, SIZE));
+    hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
+    HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
+
+    for (unsigned i = 0; i < LEN; i++) {
+        assert(i == A[i]);
+    }
+  }
+};
+
+int main() {
+  TestPlacementNew().run();
+  passed();
+}