Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

hipamd/README.md

@@ -15,11 +15,14 @@ New projects can be developed directly in the portable HIP C++ language and can
 cd HIP-privatestaging
 mkdir build
 cd build
-cmake ..
+cmake -DHSA_DIR=/path/to/hsa -DHCC_DIR=/path/to/hcc -DHIP_INSTALL_DIR=/where/to/install/hip -DCMAKE_BUILD_TYPE=Release ..
 make
-sudo make install
+make install
 ```
-Make sure HIP_PATH is pointed to `/opt/hip` and PATH includes `$HIP_PATH/bin`
+Make sure HIP_PATH is pointed to `/where/to/install/hip` and PATH includes `$HIP_PATH/bin`. This requirement is optional, but required to run any HIP test infrastructure.
+
+The Release build installs HIP inside `/where/to/install/hip`. The Debug build will install inside HIP repo directory by adding a `./lib` directory containing `libhip_hcc.a`
+
 
 ## More Info:
 - [HIP FAQ](docs/markdown/hip_faq.md)

hipamd/include/hcc_detail/hip_hcc.h

@@ -219,9 +219,9 @@ extern "C" {
 #endif
 
 typedef class ihipStream_t* hipStream_t;
-typedef struct hipEvent_t {
-    struct ihipEvent_t *_handle;
-} hipEvent_t;
+//typedef struct hipEvent_t {
+//    struct ihipEvent_t *_handle;
+//} hipEvent_t;
 
 #ifdef __cplusplus
 }

hipamd/include/hcc_detail/hip_runtime_api.h

@@ -31,7 +31,7 @@ THE SOFTWARE.
 
 #include <hcc_detail/host_defines.h>
 #include <hip_runtime_api.h>
-#include "hip_hcc.h"
+//#include "hip_hcc.h"
 
 #if defined (__HCC__) &&  (__hcc_workweek__ < 16074)
 #error("This version of HIP requires a newer version of HCC.");
@@ -42,6 +42,12 @@ THE SOFTWARE.
 extern "C" {
 #endif
 
+typedef struct ihipStream_t *hipStream_t;
+typedef struct hipEvent_t {
+    struct ihipEvent_t *_handle;
+} hipEvent_t;
+
+
 /**
  * @addtogroup GlobalDefs More
  * @{

hipamd/samples/0_Intro/bit_extract/Makefile

@@ -1,6 +1,6 @@
 #Dependencies : [MYHIP]/bin must be in user's path.
 
-HIP_PATH?=$(shell hipconfig --path)
+HIP_PATH=../../..
 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
 HIPCC=$(HIP_PATH)/bin/hipcc
 

hipamd/samples/0_Intro/square/Makefile

@@ -1,4 +1,4 @@
-HIP_PATH?=$(shell hipconfig --path)
+HIP_PATH=../../..
 HIPCC=$(HIP_PATH)/bin/hipcc
 
 all: square.hip.out

hipamd/samples/1_Utils/hipBusBandwidth/Makefile

@@ -1,4 +1,4 @@
-HIP_PATH?=$(shell hipconfig -p)
+HIP_PATH=../../..
 HIPCC=$(HIP_PATH)/bin/hipcc
 
 EXE=hipBusBandwidth

hipamd/samples/1_Utils/hipDispatchLatency/Makefile

@@ -1,4 +1,4 @@
-HIP_PATH?=$(shell hipconfig -p)
+HIP_PATH=../../..
 HIPCC=$(HIP_PATH)/bin/hipcc
 
 EXE=hipDispatchLatency

hipamd/samples/1_Utils/hipInfo/Makefile

@@ -1,4 +1,4 @@
-HIP_PATH?=$(shell hipconfig -p)
+HIP_PATH=../../..
 HIPCC=$(HIP_PATH)/bin/hipcc
 
 EXE=hipInfo

hipamd/tests/src/hipC.cpp

@@ -0,0 +1,22 @@
+#include"hip_runtime.h"
+#include<stdio.h>
+
+#define ITER 1<<20
+#define SIZE 1024*1024*sizeof(int)
+
+__global__ void Iter(hipLaunchParm lp, int *Ad){
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    if(tx == 0){
+        for(int i=0;i<ITER;i++){
+            Ad[tx] += 1;
+        }
+    }
+}
+
+int main(){
+    int A=0, *Ad;
+    hipMalloc((void**)&Ad, SIZE);
+    hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
+    hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, 0, Ad);
+    hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
+}