From 153d5f8e1b92bd01bc5b1f9e1fdc8aa602e471b7 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Sat, 9 Apr 2016 14:40:17 -0500 Subject: [PATCH] xcompilation tests --- tests/src/experimental/xcompile/gApi.c | 8 +++ tests/src/experimental/xcompile/gHipApi.c | 45 ++++++++++++++ tests/src/experimental/xcompile/gHipApi.h | 20 +++++++ tests/src/experimental/xcompile/gapi.sh | 9 +++ tests/src/experimental/xcompile/ghipapi.sh | 8 +++ tests/src/experimental/xcompile/gxxApi.cpp | 9 +++ tests/src/experimental/xcompile/gxxApi1.cpp | 9 +++ tests/src/experimental/xcompile/gxxApi1.h | 8 +++ tests/src/experimental/xcompile/gxxHipApi.cpp | 18 ++++++ tests/src/experimental/xcompile/gxxHipApi.h | 58 +++++++++++++++++++ tests/src/experimental/xcompile/gxxapi.sh | 8 +++ tests/src/experimental/xcompile/gxxhipapi.sh | 8 +++ tests/src/experimental/xcompile/hHip.c | 37 ++++++++++++ tests/src/experimental/xcompile/hHipApi.c | 21 +++++++ tests/src/experimental/xcompile/hipapig.sh | 3 + tests/src/experimental/xcompile/hipapigxx.sh | 3 + tests/src/experimental/xcompile/hipg.sh | 8 +++ tests/src/experimental/xcompile/hipgapi.sh | 8 +++ tests/src/experimental/xcompile/hipgxx.sh | 7 +++ tests/src/experimental/xcompile/hipgxxapi.sh | 7 +++ tests/src/experimental/xcompile/hipxxKer.cpp | 29 ++++++++++ tests/src/experimental/xcompile/hxxHip.cpp | 43 ++++++++++++++ tests/src/experimental/xcompile/hxxHipApi.cpp | 43 ++++++++++++++ 23 files changed, 417 insertions(+) create mode 100644 tests/src/experimental/xcompile/gApi.c create mode 100644 tests/src/experimental/xcompile/gHipApi.c create mode 100644 tests/src/experimental/xcompile/gHipApi.h create mode 100755 tests/src/experimental/xcompile/gapi.sh create mode 100755 tests/src/experimental/xcompile/ghipapi.sh create mode 100644 tests/src/experimental/xcompile/gxxApi.cpp create mode 100644 tests/src/experimental/xcompile/gxxApi1.cpp create mode 100644 tests/src/experimental/xcompile/gxxApi1.h create mode 100644 tests/src/experimental/xcompile/gxxHipApi.cpp create mode 100644 tests/src/experimental/xcompile/gxxHipApi.h create mode 100755 tests/src/experimental/xcompile/gxxapi.sh create mode 100755 tests/src/experimental/xcompile/gxxhipapi.sh create mode 100644 tests/src/experimental/xcompile/hHip.c create mode 100644 tests/src/experimental/xcompile/hHipApi.c create mode 100755 tests/src/experimental/xcompile/hipapig.sh create mode 100755 tests/src/experimental/xcompile/hipapigxx.sh create mode 100755 tests/src/experimental/xcompile/hipg.sh create mode 100755 tests/src/experimental/xcompile/hipgapi.sh create mode 100755 tests/src/experimental/xcompile/hipgxx.sh create mode 100755 tests/src/experimental/xcompile/hipgxxapi.sh create mode 100644 tests/src/experimental/xcompile/hipxxKer.cpp create mode 100644 tests/src/experimental/xcompile/hxxHip.cpp create mode 100644 tests/src/experimental/xcompile/hxxHipApi.cpp diff --git a/tests/src/experimental/xcompile/gApi.c b/tests/src/experimental/xcompile/gApi.c new file mode 100644 index 0000000000..cd3dd0717a --- /dev/null +++ b/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); +} diff --git a/tests/src/experimental/xcompile/gHipApi.c b/tests/src/experimental/xcompile/gHipApi.c new file mode 100644 index 0000000000..3a1a9c683f --- /dev/null +++ b/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; +} + + diff --git a/tests/src/experimental/xcompile/gHipApi.h b/tests/src/experimental/xcompile/gHipApi.h new file mode 100644 index 0000000000..f7bb1d8d08 --- /dev/null +++ b/tests/src/experimental/xcompile/gHipApi.h @@ -0,0 +1,20 @@ +#ifndef GHIPAPI_H +#define GHIPAPI_H + +#include + +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 diff --git a/tests/src/experimental/xcompile/gapi.sh b/tests/src/experimental/xcompile/gapi.sh new file mode 100755 index 0000000000..489171d65a --- /dev/null +++ b/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 + + diff --git a/tests/src/experimental/xcompile/ghipapi.sh b/tests/src/experimental/xcompile/ghipapi.sh new file mode 100755 index 0000000000..611053896d --- /dev/null +++ b/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 + diff --git a/tests/src/experimental/xcompile/gxxApi.cpp b/tests/src/experimental/xcompile/gxxApi.cpp new file mode 100644 index 0000000000..21488028ef --- /dev/null +++ b/tests/src/experimental/xcompile/gxxApi.cpp @@ -0,0 +1,9 @@ +#include"hip_runtime_api.h" +#include + +#define size 1024*1024 + +int main(){ + float *Ad; + hipMalloc((void**)&Ad, size); +} diff --git a/tests/src/experimental/xcompile/gxxApi1.cpp b/tests/src/experimental/xcompile/gxxApi1.cpp new file mode 100644 index 0000000000..6a76fc06bc --- /dev/null +++ b/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; +} diff --git a/tests/src/experimental/xcompile/gxxApi1.h b/tests/src/experimental/xcompile/gxxApi1.h new file mode 100644 index 0000000000..74d04d4549 --- /dev/null +++ b/tests/src/experimental/xcompile/gxxApi1.h @@ -0,0 +1,8 @@ +#ifndef GXXAPI1_H +#define GXXAPI1_H +#include + +void* mallocHip(size_t size); + + +#endif diff --git a/tests/src/experimental/xcompile/gxxHipApi.cpp b/tests/src/experimental/xcompile/gxxHipApi.cpp new file mode 100644 index 0000000000..99d496f6aa --- /dev/null +++ b/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); +} diff --git a/tests/src/experimental/xcompile/gxxHipApi.h b/tests/src/experimental/xcompile/gxxHipApi.h new file mode 100644 index 0000000000..5409100c9c --- /dev/null +++ b/tests/src/experimental/xcompile/gxxHipApi.h @@ -0,0 +1,58 @@ +#ifndef GXXHIPAPI_H +#define GXXHIPAPI_H + +#include +#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 + void setDevPtr(T* ptr) + { + devPtr = (void*)ptr; + } + + template + T* getDevPtr() + { + return (T*)devPtr; + } + + template + void setHstPtr(T* ptr) + { + hstPtr = (void*)ptr; + } + + template + T* getHstPtr() + { + return (T*)hstPtr; + } + + void H2D(); + void D2H(); + template + void hostMemSet(T val) + { + T* tmpPtr = (T*)hstPtr; + for(int i=0;i + void memAlloc() + { + hipMalloc((void**)&devPtr, size); + } +}; + +#endif diff --git a/tests/src/experimental/xcompile/gxxapi.sh b/tests/src/experimental/xcompile/gxxapi.sh new file mode 100755 index 0000000000..42569f93d2 --- /dev/null +++ b/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 + diff --git a/tests/src/experimental/xcompile/gxxhipapi.sh b/tests/src/experimental/xcompile/gxxhipapi.sh new file mode 100755 index 0000000000..2405b15cd1 --- /dev/null +++ b/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 + diff --git a/tests/src/experimental/xcompile/hHip.c b/tests/src/experimental/xcompile/hHip.c new file mode 100644 index 0000000000..3e4f5e3041 --- /dev/null +++ b/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); + + +} diff --git a/tests/src/experimental/xcompile/hHipApi.c b/tests/src/experimental/xcompile/hHipApi.c new file mode 100644 index 0000000000..7761d656c1 --- /dev/null +++ b/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); +} + + + diff --git a/tests/src/experimental/xcompile/hipapig.sh b/tests/src/experimental/xcompile/hipapig.sh new file mode 100755 index 0000000000..c0f96d6180 --- /dev/null +++ b/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 diff --git a/tests/src/experimental/xcompile/hipapigxx.sh b/tests/src/experimental/xcompile/hipapigxx.sh new file mode 100755 index 0000000000..832cc878a0 --- /dev/null +++ b/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 diff --git a/tests/src/experimental/xcompile/hipg.sh b/tests/src/experimental/xcompile/hipg.sh new file mode 100755 index 0000000000..b0b6333874 --- /dev/null +++ b/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 + diff --git a/tests/src/experimental/xcompile/hipgapi.sh b/tests/src/experimental/xcompile/hipgapi.sh new file mode 100755 index 0000000000..3804018719 --- /dev/null +++ b/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 + diff --git a/tests/src/experimental/xcompile/hipgxx.sh b/tests/src/experimental/xcompile/hipgxx.sh new file mode 100755 index 0000000000..58b09c2ce9 --- /dev/null +++ b/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 diff --git a/tests/src/experimental/xcompile/hipgxxapi.sh b/tests/src/experimental/xcompile/hipgxxapi.sh new file mode 100755 index 0000000000..fd04425c47 --- /dev/null +++ b/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 diff --git a/tests/src/experimental/xcompile/hipxxKer.cpp b/tests/src/experimental/xcompile/hipxxKer.cpp new file mode 100644 index 0000000000..3e5294e62a --- /dev/null +++ b/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 +#include"hip_runtime.h" + +#define LEN 1024*1024 +#define SIZE LEN * sizeof(float) + +class memManager; + +template +__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 Vec(3); + for(int i=0;i(); + } + + for(int i=0;i(), Vec[1].getDevPtr(), Vec[2].getDevPtr(), LEN); + + Vec[2].D2H(); + assert(Vec[0].getHstPtr()[10] + Vec[1].getHstPtr()[10] == Vec[2].getHstPtr()[10]); +} diff --git a/tests/src/experimental/xcompile/hxxHipApi.cpp b/tests/src/experimental/xcompile/hxxHipApi.cpp new file mode 100644 index 0000000000..98912d4041 --- /dev/null +++ b/tests/src/experimental/xcompile/hxxHipApi.cpp @@ -0,0 +1,43 @@ +#include"gxxHipApi.h" +#include +#include"hip_runtime.h" + +#define LEN 1024*1024 +#define SIZE LEN * sizeof(float) + +class memManager; + +int main() +{ + std::vector Vec(4); + for(int i=0;i(); + } + + for(int i=0;i<2;i++) + { + Vec[i].hostMemSet((i+1)*1.0f); + Vec[i].H2D(); + } + + Vec[2].setDevPtr(Vec[0].getDevPtr()); + Vec[3].setDevPtr(Vec[1].getDevPtr()); + + for(int i=2;i()[10] == Vec[2].getHstPtr()[10]); + assert(Vec[1].getHstPtr()[10] == Vec[3].getHstPtr()[10]); +}