Added hipHostAlloc feature for CUDA

[ROCm/hip commit: f479531be5]
2016-03-05 13:58:56 -06:00
parent 5e9d9cbabf
commit 8e0fc269d7
4 changed files with 63 additions and 3 deletions
@@ -0,0 +1,42 @@
+#include"test_common.h"
+
+#define LEN 1024*1024
+#define SIZE LEN*sizeof(float)
+
+__global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){
+int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+Cd[tx] = Ad[tx] + Bd[tx];
+}
+
+int main(){
+float *A, *B, *C;
+float *Ad, *Bd, *Cd;
+
+hipDeviceProp_t prop;
+int device;
+HIPCHECK(hipGetDevice(&device));
+HIPCHECK(hipDeviceGetProperties(&prop, device));
+if(prop.canMapHostMemory != 1){
+std::cout<<"Exiting..."<<std::endl;
+}
+HIPCHECK(hipHostAlloc((void**)&A, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
+HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
+HIPCHECK(hipHostAlloc((void**)&C, SIZE, hipHostAllocMapped));
+
+HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Bd, B, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
+
+for(int i=0;i<LEN;i++){
+A[i] = 1.0f;
+B[i] = 2.0f;
+}
+
+dim3 dimGrid(LEN/512,1,1);
+dim3 dimBlock(512,1,1);
+
+hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
+
+passed();
+
+}