xcompilation tests

tests/src/experimental/xcompile/gApi.c

@@ -0,0 +1,8 @@
+#include"hip_runtime_api.h"
+
+#define size 1024*1024
+
+int main(){
+	float *Ad;
+	hipMalloc((void**)&Ad, size);
+}

tests/src/experimental/xcompile/gHipApi.c

@@ -0,0 +1,45 @@
+#include"gHipApi.h"
+#include"hip_runtime_api.h"
+#include"stdio.h"
+
+void _h2d(mem_manager *self){
+    hipMemcpy(self->dev_ptr, self->hst_ptr, self->size, hipMemcpyHostToDevice);
+}
+
+void _d2h(mem_manager *self)
+{
+    hipMemcpy(self->hst_ptr, self->dev_ptr, self->size, hipMemcpyDeviceToHost);
+}
+
+void _malloc_hip(mem_manager *self)
+{
+      hipMalloc(&(self->dev_ptr), self->size);
+}
+
+void _malloc_hst(mem_manager *self)
+{
+     self->hst_ptr = malloc(self->size);
+}
+
+void memset_hst(mem_manager *mem, float val)
+{
+      float *tmp = (float*)mem->hst_ptr;
+      int i;
+      for(i=0;i<(mem->size)/sizeof(float);i++)
+      {
+          tmp[i] = val;
+      }
+}
+
+mem_manager *mem_manager_start(size_t _size)
+{
+    mem_manager *tmp = (mem_manager*)malloc(sizeof(mem_manager));
+    tmp->size = _size;
+    tmp->h2d = _h2d;
+    tmp->d2h = _d2h;
+    tmp->malloc_hip = _malloc_hip;
+    tmp->malloc_hst = _malloc_hst;
+    return tmp;
+}
+
+

tests/src/experimental/xcompile/gHipApi.h

@@ -0,0 +1,20 @@
+#ifndef GHIPAPI_H
+#define GHIPAPI_H
+
+#include<stdlib.h>
+
+typedef struct {
+void *hst_ptr;
+void *dev_ptr;
+size_t size;
+void (*h2d)();
+void (*d2h)();
+void (*malloc_hip)();
+void (*malloc_hst)();
+} mem_manager;
+
+mem_manager *mem_manager_start(size_t);
+
+void memset_hst(mem_manager*, float);
+
+#endif

tests/src/experimental/xcompile/gapi.sh

@@ -0,0 +1,9 @@
+#!/bin/bash
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include hHipApi.o gHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi
+
+

tests/src/experimental/xcompile/ghipapi.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o
+
+hipcc hHipApi.o gHipApi.o -o gHipApi
+

tests/src/experimental/xcompile/gxxApi.cpp

@@ -0,0 +1,9 @@
+#include"hip_runtime_api.h"
+#include<iostream>
+
+#define size 1024*1024
+
+int main(){
+	float *Ad;
+	hipMalloc((void**)&Ad, size);
+}

tests/src/experimental/xcompile/gxxApi1.cpp

@@ -0,0 +1,9 @@
+#include"gxxApi1.h"
+#include"hip_runtime_api.h"
+
+void* mallocHip(size_t size)
+{
+	void *ptr;
+	hipMalloc(&ptr, size);
+	return ptr;
+}

tests/src/experimental/xcompile/gxxApi1.h

@@ -0,0 +1,8 @@
+#ifndef GXXAPI1_H
+#define GXXAPI1_H
+#include<stdlib.h>
+
+void* mallocHip(size_t size);
+
+
+#endif

tests/src/experimental/xcompile/gxxHipApi.cpp

@@ -0,0 +1,18 @@
+#include"gxxHipApi.h"
+
+memManager::memManager(const memManager &obj)
+{
+    devPtr = obj.devPtr;
+    hstPtr = obj.hstPtr;
+    size = obj.size;
+}
+
+void memManager::H2D()
+{
+    hipMemcpy(devPtr, hstPtr, size, hipMemcpyHostToDevice);
+}
+
+void memManager::D2H()
+{
+    hipMemcpy(hstPtr, devPtr, size, hipMemcpyDeviceToHost);
+}

tests/src/experimental/xcompile/gxxHipApi.h

@@ -0,0 +1,58 @@
+#ifndef GXXHIPAPI_H
+#define GXXHIPAPI_H
+
+#include<stdlib.h>
+#include"hip_runtime_api.h"
+
+class memManager{
+private:
+    void* devPtr;
+    void* hstPtr;
+    size_t size;
+public:
+    memManager(size_t size) : size(size) {}
+    memManager(){}
+    memManager(const memManager &obj);
+    template<typename T>
+    void setDevPtr(T* ptr)
+    {
+        devPtr = (void*)ptr;
+    }
+
+    template<typename T>
+    T* getDevPtr()
+    {
+         return (T*)devPtr;
+    }
+
+    template<typename T>
+    void setHstPtr(T* ptr)
+    {
+         hstPtr = (void*)ptr;
+    }
+
+    template<typename T>
+    T* getHstPtr()
+    {
+         return (T*)hstPtr;
+    }
+
+    void H2D();
+    void D2H();
+    template<typename T>
+    void hostMemSet(T val)
+    {
+        T* tmpPtr = (T*)hstPtr;
+        for(int i=0;i<size/sizeof(T);i++)
+        {
+            tmpPtr[i] = val;
+        }
+    }
+    template<typename T>
+    void memAlloc()
+    {
+          hipMalloc((void**)&devPtr, size);
+    }
+};
+
+#endif

tests/src/experimental/xcompile/gxxapi.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hxxHipApi.o
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi
+

tests/src/experimental/xcompile/gxxhipapi.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+
+hipcc -c hxxHipApi.cpp -o hxxHipApi.o
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi
+

tests/src/experimental/xcompile/hHip.c

@@ -0,0 +1,37 @@
+#include "gHipApi.h"
+#include "hip_runtime.h"
+
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(float)
+
+__global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd, size_t len)
+{
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    if(tx < len)
+    {
+        Cd[tx] = Ad[tx] + Bd[tx];
+    }
+}
+
+int main()
+{
+    mem_manager *a, *b, *c;
+    a = mem_manager_start(SIZE);
+    b = mem_manager_start(SIZE);
+    c = mem_manager_start(SIZE);
+    a->malloc_hst(a);
+    b->malloc_hst(b);
+    c->malloc_hst(c);
+    a->malloc_hip(a);
+    b->malloc_hip(b);
+    c->malloc_hip(c);
+    memset_hst(a, 1.0f);
+    memset_hst(b, 2.0f);
+    a->h2d(a);
+    b->h2d(b);
+    hipLaunchKernel(HIP_KERNEL_NAME(Add), dim3(LEN/1024), dim3(1024), 0, 0, (float*)a->dev_ptr, (float*)b->dev_ptr, (float*)c->dev_ptr, LEN);
+    c->d2h(c);
+    assert(((float*)c->hst_ptr)[10] == 3.0f);
+
+
+}

tests/src/experimental/xcompile/hHipApi.c

@@ -0,0 +1,21 @@
+#include"gHipApi.h"
+#include"stdio.h"
+#include "assert.h"
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(float)
+
+int main()
+{
+    mem_manager *a;
+    a = mem_manager_start(SIZE);
+    a->malloc_hst(a);
+    a->malloc_hip(a);
+    memset_hst(a, 1.0f);
+    a->h2d(a);
+    memset_hst(a, 0.0f);
+    a->d2h(a);
+    assert(((float*)a->hst_ptr)[10] == 1.0f);
+}
+
+
+

tests/src/experimental/xcompile/hipapig.sh

@@ -0,0 +1,3 @@
+#!/bin/bash
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi

tests/src/experimental/xcompile/hipapigxx.sh

@@ -0,0 +1,3 @@
+#!/bin/bash
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi

tests/src/experimental/xcompile/hipg.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+
+hipcc -c hHip.c -o hHip.o
+
+hipcc hHip.o gHipApi.o -o hipG
+

tests/src/experimental/xcompile/hipgapi.sh

@@ -0,0 +1,8 @@
+#!/bin/bash
+
+gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o
+
+hipcc -c hHipApi.c -o hHipApi.o
+
+hipcc hHipApi.o gHipApi.o -o hipGApi
+

tests/src/experimental/xcompile/hipgxx.sh

@@ -0,0 +1,7 @@
+#!/bin/bash
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+
+hipcc -c hxxHip.cpp -o hxxHip.o
+
+hipcc hxxHip.o gxxHipApi.o -o hxxHip

tests/src/experimental/xcompile/hipgxxapi.sh

@@ -0,0 +1,7 @@
+#!/bin/bash
+
+g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o
+
+hipcc -c hxxHipApi.cpp -o hxxHipApi.o
+
+hipcc hxxHipApi.o gxxHipApi.o -o hxxHipApi

tests/src/experimental/xcompile/hipxxKer.cpp

@@ -0,0 +1,29 @@
+#include"hip_runtime.h"
+#include"hip_runtime_api.h"
+#include"gxxApi1.h"
+
+#define len 1024*1024
+#define size len * sizeof(float)
+
+__global__ void Kern(hipLaunchParm lp, float *A)
+{
+	int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+	A[tx] += 1.0f;
+}
+
+int main()
+{
+	float *A, *Ad;
+	for(int i=0;i<len;i++)
+	{
+		A[i] = 1.0f;
+	}
+	Ad = (float*)mallocHip(size);
+	memcpyHipH2D(Ad, A, size);
+	hipLaunchKernel(HIP_KERNEL_NAME(Kern), dim3(len/1024), dim3(1024), 0, 0, A);
+	memcpyHipD2H(A, Ad, size);
+	for(int i=0;i<len;i++)
+	{
+		assert(A[i] == 2.0f);
+	}
+}

tests/src/experimental/xcompile/hxxHip.cpp

@@ -0,0 +1,43 @@
+#include"gxxHipApi.h"
+#include<vector>
+#include"hip_runtime.h"
+
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(float)
+
+class memManager;
+
+template<typename T>
+__global__ void Add(hipLaunchParm lp, T* Ad, T* Bd, T* Cd, size_t Len)
+{
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    if(tx < Len)
+    {
+        Cd[tx] = Ad[tx] + Bd[tx];
+    }
+}
+
+int main()
+{
+    std::vector<class memManager> Vec(3);
+    for(int i=0;i<Vec.size();i++){
+        Vec[i] = memManager(SIZE);
+    }
+
+    for(int i=0;i<3;i++)
+    {
+        Vec[i].setHstPtr(new float[LEN]);
+        Vec[i].memAlloc<float>();
+    }
+
+    for(int i=0;i<Vec.size()-1;i++)
+    {
+        Vec[i].hostMemSet((i+1)*1.0f);
+        Vec[i].H2D();
+    }
+
+    hipLaunchKernel(HIP_KERNEL_NAME(Add), dim3(LEN/1024), dim3(1024), 0, 0, Vec[0].getDevPtr<float>(), Vec[1].getDevPtr<float>(), Vec[2].getDevPtr<float>(), LEN);
+
+    Vec[2].D2H();
+    assert(Vec[0].getHstPtr<float>()[10] + Vec[1].getHstPtr<float>()[10] == Vec[2].getHstPtr<float>()[10]);
+}

tests/src/experimental/xcompile/hxxHipApi.cpp

@@ -0,0 +1,43 @@
+#include"gxxHipApi.h"
+#include<vector>
+#include"hip_runtime.h"
+
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(float)
+
+class memManager;
+
+int main()
+{
+    std::vector<class memManager> Vec(4);
+    for(int i=0;i<Vec.size();i++){
+        Vec[i] = memManager(SIZE);
+    }
+
+    for(int i=0;i<4;i++)
+    {
+        Vec[i].setHstPtr(new float[LEN]);
+    }
+
+    for(int i=0;i<2;i++)
+    {
+        Vec[i].memAlloc<float>();
+    }
+
+    for(int i=0;i<2;i++)
+    {
+        Vec[i].hostMemSet((i+1)*1.0f);
+        Vec[i].H2D();
+    }
+
+    Vec[2].setDevPtr(Vec[0].getDevPtr<float>());
+    Vec[3].setDevPtr(Vec[1].getDevPtr<float>());
+
+    for(int i=2;i<Vec.size();i++)
+    {
+        Vec[i].D2H();
+    }
+    
+    assert(Vec[0].getHstPtr<float>()[10] == Vec[2].getHstPtr<float>()[10]);
+    assert(Vec[1].getHstPtr<float>()[10] == Vec[3].getHstPtr<float>()[10]);
+}