src/ fixed hipHostAllocDefault flags

[ROCm/hip commit: b768568902]

projects/hip/src/hip_hcc.cpp

@@ -2054,31 +2054,31 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){
 	auto device = ihipGetTlsDefaultDevice();
 
 	if(device){
-		if(flags & hipHostAllocDefault){
-		const unsigned am_flags = amHostPinned;
+		if(flags == hipHostAllocDefault){
+		    const unsigned am_flags = amHostPinned;
 
-		*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
-		if(sizeBytes && (*ptr == NULL)){
-			hip_status = hipErrorMemoryAllocation;
-		}else{
-			hc::am_memtracker_update(*ptr, device->_device_index, 0);
-		}
-		tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
+		    *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
+		    if(sizeBytes && (*ptr == NULL)){
+			    hip_status = hipErrorMemoryAllocation;
+		    }else{
+			    hc::am_memtracker_update(*ptr, device->_device_index, 0);
+		    }
+		    tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
 		}
 		if(flags & hipHostAllocMapped){
-		const unsigned am_flags = amHostPinned;
+		    const unsigned am_flags = amHostPinned;
 
-		*ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
-		if(sizeBytes && (*ptr == NULL)){
-			hip_status = hipErrorMemoryAllocation;
-		}else{
-			hc::am_memtracker_update(*ptr, device->_device_index, flags);
+		    *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
+		    if(sizeBytes && (*ptr == NULL)){
+			    hip_status = hipErrorMemoryAllocation;
+		    }else{
+			    hc::am_memtracker_update(*ptr, device->_device_index, flags);
 //			void *srcPtr;
 //			hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr);
 //			assert(hsa_status == HSA_STATUS_SUCCESS);
 //			hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false);
-		}
-		tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
+		    }
+		    tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
 		}
 	}
 	return ihipLogStatus(hip_status);

projects/hip/tests/src/hipHostAlloc.cpp

@@ -39,14 +39,14 @@ int device;
 HIPCHECK(hipGetDevice(&device));
 HIPCHECK(hipGetDeviceProperties(&prop, device));
 if(prop.canMapHostMemory != 1){
-std::cout<<"Exiting..."<<std::endl;
+//std::cout<<"Exiting..."<<std::endl;
 }
+
 HIPCHECK(hipHostAlloc((void**)&A, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
-HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
+HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocDefault));
 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++){
@@ -54,6 +54,9 @@ A[i] = 1.0f;
 B[i] = 2.0f;
 }
 
+HIPCHECK(hipMalloc((void**)&Bd, SIZE));
+HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+
 dim3 dimGrid(LEN/512,1,1);
 dim3 dimBlock(512,1,1);