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

Conflicts:
	src/hip_hcc.cpp
	tests/src/CMakeLists.txt

RELEASE.md

@@ -23,7 +23,7 @@ Date: 2016.02.18
 - Update Runtime Documentation.
 - Improve implementations of cross-lane operations (_ballot, _any, _all).
 - Provide shuffle intrinsics (performance optimization in-progress).
-- Support hipDeviceAttribute for querying "one-shot" device attributes, as an alternative to hipDeviceGetProperties.
+- Support hipDeviceAttribute for querying "one-shot" device attributes, as an alternative to hipGetDeviceProperties.
 
 
 ===================================================================================================

bin/hipify

@@ -286,7 +286,12 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaMallocHost\b/hipMallocHost/g;
         $ft{'mem'} += s/\bcudaFree\b/hipFree/g;
         $ft{'mem'} += s/\bcudaFreeHost\b/hipFreeHost/g;
-
+		$ft{'mem'} += s/\bcudaHostAlloc\b/hipHostAlloc/g;
+		$ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
+		$ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostAllocDefault/g;
+		$ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostAllocPortable/g;
+		$ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostAllocMapped/g;
+		$ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostAllocWriteCombined/g;
 
 
         #--------
@@ -345,7 +350,7 @@ while (@ARGV) {
         #--------
         # Device 
         $ft{'dev'} += s/\bcudaDeviceProp\b/hipDeviceProp_t/g;
-        $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipDeviceGetProperties/g;
+        $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipGetDeviceProperties/g;
 
         # Attribute
         $ft{'err'} += s/\bcudaDevAttrMaxThreadsPerBlock\b/hipDeviceAttributeMaxThreadsPerBlock/g;

docs/markdown/hip_kernel_language.md

@@ -427,7 +427,7 @@ Nvidia devices implement the timer as a per-compute-unit clock that increments o
 To obtain the clock frequency, use the hipDeviceProp_t.clockInstructionRate field:
 
 ```
-hipDeviceGetProperties(&deviceProps, deviceId);
+hipGetDeviceProperties(&deviceProps, deviceId);
 // Compute time in ms--device_ticks is based on values reported from clock() device function
 float time = device_ticks / (float)deviceProps.clockInstructionRate;
 ```

docs/markdown/hip_porting_guide.md

@@ -215,10 +215,10 @@ For host code, the `__HIP_ARCH__*` defines are set to 0. You should only use the
 
 ### Device-Architecture Properties
 
-Host code should query the architecture feature flags in the device properties that hipDeviceGetProperties returns, rather than testing the "major" and "minor" fields directly:
+Host code should query the architecture feature flags in the device properties that hipGetDeviceProperties returns, rather than testing the "major" and "minor" fields directly:
 
 ```
-hipDeviceGetProperties(&deviceProp, device);
+hipGetDeviceProperties(&deviceProp, device);
 //if ((deviceProp.major == 1 && deviceProp.minor < 2))  // non-portable
 if (deviceProp.arch.hasSharedInt32Atomics) {            // portable HIP feature query
     // has shared int32 atomic operations ...

include/hcc_detail/hip_runtime_api.h

@@ -56,6 +56,16 @@ extern "C" {
 #define hipEventInterprocess  0x4  ///< Event can support IPC.  @warning - not supported in HIP.
 
 
+#define hipHostAllocDefault 0x0
+#define hipHostAllocPortable 0x1
+#define hipHostAllocMapped 0x2
+#define hipHostAllocWriteCombined 0x4
+
+#define hipHostRegisterDefault 0x0
+#define hipHostRegisterPortable 0x1 
+#define hipHostRegisterMapped 0x2
+#define hipHostRegisterIoMemory 0x4
+
 /**
  * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored.
  */
@@ -247,9 +257,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
  * @param [out] prop written with device properties
  * @param [in]  device which device to query for information
  *
- * Populates hipDeviceGetProperties with information for the specified device.
+ * Populates hipGetDeviceProperties with information for the specified device.
  */
-hipError_t hipDeviceGetProperties(hipDeviceProp_t* prop, int device);
+hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int device);
 
 
 
@@ -662,6 +672,53 @@ hipError_t hipMalloc(void** ptr, size_t size) ;
  */
 hipError_t hipMallocHost(void** ptr, size_t size) ;
 
+/**
+ *  @brief Allocate device accessible page locked host memory 
+ *
+ *  @param[out]  ptr Pointer to the allocated host pinned memory
+ *  @param[in] size Requested memory size
+ *  @param[in] flags Type of host memory allocation
+ *  @return Error code
+ */
+hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) ;
+
+/**
+ *  @brief Get Device pointer from Host Pointer allocated through hipHostAlloc
+ *
+ *  @param[out]  dstPtr Device Pointer mapped to passed host pointer
+ *  @param[in] hstPtr Host Pointer allocated through hipHostAlloc
+ *  @param[in] size Requested memory size
+ *  @return Error code
+ */
+hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size) ;
+
+/**
+ *  @brief Get flags associated with host pointer
+ *
+ *  @param[out]  flagsPtr Memory location to store flags
+ *  @param[in] hostPtr Host Pointer allocated through hipHostAlloc
+ *  @return Error code
+ */
+hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ;
+
+/**
+ *  @brief Pin host memory
+ *
+ *  @param[out] hostPtr Pointer to host memory to be pinned
+ *  @param[in] sizeBytes size of the host memory
+ *  @param[in] flags Type of pinning the the host memory
+ *  @return Error code
+ */
+hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) ;
+
+/**
+ *  @brief Un-pin host pointer
+ *
+ *  @param[in] hostPtr Pinned Host Pointer
+ *  @return Error code
+ */
+hipError_t hipHostUnregister(void* hostPtr) ;
+
 
 /**
  *  @brief Free memory allocated by the hcc hip memory allocation API.

include/hip_runtime_api.h

@@ -97,6 +97,7 @@ typedef struct hipDeviceProp_t {
     int pciDeviceID;                            ///< PCI Device ID.
     size_t maxSharedMemoryPerMultiProcessor;    ///< Maximum Shared Memory Per Multiprocessor.
     int isMultiGpuBoard;                        ///< 1 if device is on a multi-GPU board, 0 if not.
+    int canMapHostMemory;                       ///< Check whether HIP can map host memory
  } hipDeviceProp_t;
 
 

include/nvcc_detail/hip_runtime_api.h

@@ -50,6 +50,13 @@ hipMemcpyHostToHost
 } hipTextureFilterMode;*/
 #define hipFilterModePoint cudaFilterModePoint
 
+#define hipHostAllocDefault cudaHostAllocDefault
+#define hipHostAllocPortable cudaHostAllocPortable
+#define hipHostAllocMapped cudaHostAllocMapped
+#define hipHostAllocWriteCombined cudaHostAllocWriteCombined
+
+#define hipHostRegisterPortable cudaHostRegisterPortable
+#define hipHostRegisterMapped cudaHostRegisterMapped
 
 typedef cudaEvent_t hipEvent_t;
 typedef cudaStream_t hipStream_t;
@@ -115,6 +122,27 @@ inline static hipError_t hipFree(void* ptr) {
 inline static hipError_t hipMallocHost(void** ptr, size_t size) {
     return hipCUDAErrorTohipError(cudaMallocHost(ptr, size));
 }
+
+inline static hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
+}
+
+inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags));
+}
+
+inline static hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr){
+	return hipCUDAErrorTohipError(cudaHostGetFlags(flagsPtr, hostPtr));
+}
+
+inline static hipError_t hipHostRegister(void* ptr, size_t size, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostRegister(ptr, size, flags));
+}
+
+inline static hipError_t hipHostUnregister(void* ptr){
+	return hipCUDAErrorTohipError(cudaHostUnregister(ptr));
+}
+
 inline static hipError_t hipFreeHost(void* ptr) {
     return hipCUDAErrorTohipError(cudaFreeHost(ptr));
 }
@@ -154,7 +182,7 @@ inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
     return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
 }
 
-inline static hipError_t hipDeviceGetProperties(hipDeviceProp_t *p_prop, int device)
+inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device)
 {
 	cudaDeviceProp cdprop;
 	cudaError_t cerror;
@@ -177,6 +205,7 @@ inline static hipError_t hipDeviceGetProperties(hipDeviceProp_t *p_prop, int dev
 	p_prop->l2CacheSize = cdprop.l2CacheSize ;
 	p_prop->maxThreadsPerMultiProcessor = cdprop.maxThreadsPerMultiProcessor ;
 	p_prop->computeMode = cdprop.computeMode ;
+	p_prop->canMapHostMemory = cdprop.canMapHostMemory;
 
 	// Same as clock-rate:
 	p_prop->clockInstructionRate = cdprop.clockRate; 

samples/0_Intro/bit_extract/bit_extract.cpp

@@ -60,7 +60,7 @@ int main(int argc, char *argv[])
     int deviceId;
     CHECK (hipGetDevice(&deviceId));
 	hipDeviceProp_t props;
-	CHECK(hipDeviceGetProperties(&props, deviceId));
+	CHECK(hipGetDeviceProperties(&props, deviceId));
 	printf ("info: running on device #%d %s\n", deviceId, props.name);
 
 

samples/0_Intro/square/square.hipref.cpp

@@ -53,7 +53,7 @@ int main(int argc, char *argv[])
 	size_t Nbytes = N * sizeof(float);
 
 	hipDeviceProp_t props;
-	CHECK(hipDeviceGetProperties(&props, 0/*deviceID*/));
+	CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
 	printf ("info: running on device %s\n", props.name);
 
 	printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);

samples/1_Utils/hipInfo/hipInfo.cpp

@@ -73,7 +73,7 @@ void printDeviceProp (int deviceId)
     cout << setw(w1) << "device#" << deviceId << endl;
 
     hipDeviceProp_t props;
-    HIPCHECK(hipDeviceGetProperties(&props, deviceId));
+    HIPCHECK(hipGetDeviceProperties(&props, deviceId));
 
     cout << setw(w1) << "Name: " << props.name << endl;
     cout << setw(w1) << "pciBusID: " << props.pciBusID << endl;

src/hip_hcc.cpp

@@ -1435,7 +1435,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
  * @bug HCC always returns 0 for regsPerBlock
  * @bug HCC always returns 0 for l2CacheSize
  */
-hipError_t hipDeviceGetProperties(hipDeviceProp_t* props, int device)
+hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device)
 {
     std::call_once(hip_initialized, ihipInit);
 
@@ -2045,6 +2045,137 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
     return ihipLogStatus(hip_status);
 }
 
+
+hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){
+	std::call_once(hip_initialized, ihipInit);
+
+	hipError_t hip_status = hipSuccess;
+
+	auto device = ihipGetTlsDefaultDevice();
+
+	if(device){
+		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);
+		}
+		if(flags & hipHostAllocMapped){
+		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);
+//			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);
+		}
+	}
+	return ihipLogStatus(hip_status);
+}
+
+
+hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size){
+	std::call_once(hip_initialized, ihipInit);
+
+	hipError_t hip_status = hipSuccess;
+
+	if(hstPtr == NULL){
+		hip_status = hipErrorInvalidValue;
+	}else{
+
+	hc::accelerator acc;
+	hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
+	am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hstPtr);
+	if(status == AM_SUCCESS){
+		*devPtr = amPointerInfo._devicePointer;
+		if(devPtr == NULL){
+			hip_status = hipErrorMemoryAllocation;
+        }
+	}
+	tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *devPtr);
+	}
+	return ihipLogStatus(hip_status);
+}
+
+hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
+{
+	std::call_once(hip_initialized, ihipInit);
+	hipError_t hip_status = hipSuccess;
+
+	hc::accelerator acc;
+	hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
+	am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
+	if(status == AM_SUCCESS){
+		*flagsPtr = amPointerInfo._appAllocationFlags;
+		if(*flagsPtr == 0){
+			hip_status = hipErrorInvalidValue;
+		}
+		else{
+			hip_status = hipSuccess;
+		}
+		tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
+	}else{
+		hip_status = hipErrorInvalidValue;
+	}
+	return ihipLogStatus(hip_status);
+}
+
+hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
+{
+	std::call_once(hip_initialized, ihipInit);
+	hipError_t hip_status = hipSuccess;
+
+	auto device = ihipGetTlsDefaultDevice();
+	void* srcPtr;
+	if(hostPtr == NULL){
+		return ihipLogStatus(hipErrorInvalidValue);
+	}
+	if(device){
+	if(flags == hipHostAllocDefault){
+		hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
+		if(hsa_status == HSA_STATUS_SUCCESS){
+			hip_status = hipSuccess;
+		}else{
+			hip_status = hipErrorMemoryAllocation;
+		}
+	}
+	else if (flags | hipHostRegisterMapped){
+		hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
+		//TODO: Added feature for actual host pointer being tracked
+		if(hsa_status != HSA_STATUS_SUCCESS){
+			hip_status = hipErrorMemoryAllocation;
+		}
+	}
+	}
+	return ihipLogStatus(hip_status);
+}
+
+hipError_t hipHostUnregister(void *hostPtr){
+	std::call_once(hip_initialized, ihipInit);
+	hipError_t hip_status = hipSuccess;
+	if(hostPtr == NULL){
+		hip_status = hipErrorInvalidValue;
+	}else{
+	hsa_status_t hsa_status = hsa_amd_memory_unlock(hostPtr);
+	if(hsa_status != HSA_STATUS_SUCCESS){
+		hip_status = hipErrorInvalidValue;
+// TODO: Add a different return error. This is not true
+	}
+	}
+	return ihipLogStatus(hip_status);
+}
+
 //---
 hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
 {
@@ -2433,6 +2564,8 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
 
     stream = ihipSyncAndResolveStream(stream);
 
+	bool trueAsync = true;
+
     if (stream) {
         ihipDevice_t *device = stream->getDevice();
 
@@ -2450,28 +2583,47 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
             memcpy(dst, src, sizeBytes);
 
         } else {
+			hc::accelerator acc;
+			hc::AmPointerInfo dstAm(NULL, NULL, 0, acc, 0, 0);
+			hc::AmPointerInfo srcAm(NULL, NULL, 0, acc, 0, 0);
+			am_status_t statDst = hc::am_memtracker_getinfo(&dstAm, dst);
+			am_status_t statSrc = hc::am_memtracker_getinfo(&srcAm, src);
+
+			if(dstAm._appAllocationFlags != 1 || srcAm._appAllocationFlags != 1){
+				trueAsync = false;
+			}
+
 			if (kind == hipMemcpyDefault) {
-				std::cout<<"hipMemcpyDefault"<<std::endl;
-				hipPointerAttribute_t att;
-				hipError_t hipSt = hipPointerGetAttributes(&att, dst);
-				if(hipSt == hipSuccess){
-					if(att.devicePointer != NULL && att.hostPointer != NULL){
-						return hipSuccess;
+				if(statDst == AM_SUCCESS && statSrc == AM_SUCCESS){
+					if(dstAm._devicePointer != NULL){
+						if(srcAm._devicePointer != NULL){
+							kind = hipMemcpyDeviceToDevice;
+						}
+						if(srcAm._hostPointer != NULL){
+							kind = hipMemcpyHostToDevice;
+						}
+					}
+					if(dstAm._hostPointer != NULL){
+						if(srcAm._devicePointer != NULL){
+							kind = hipMemcpyDeviceToHost;
+						}
+						if(srcAm._hostPointer != NULL){
+							kind = hipMemcpyHostToHost;
+						}
 					}
 				}
-				hipSt = hipPointerGetAttributes(&att, (void*)src);
-				if(hipSt == hipSuccess){
-					if(att.devicePointer != NULL && att.hostPointer != NULL){
-						return hipSuccess;
-					}
+				else{
+					return hipErrorInvalidMemcpyDirection;
 				}
-				else{return hipErrorInvalidMemcpyDirection;}
 			}
             ihipSignal_t *ihip_signal = stream->getSignal();
             hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
 
             ihipCommand_t copyType;
-            if ((kind == hipMemcpyHostToDevice) || (kind == hipMemcpyDeviceToDevice)) {
+            if (kind == hipMemcpyHostToDevice ){
+				copyType = ihipCommandCopyH2D;
+				
+			}else if(kind == hipMemcpyDeviceToDevice) {
                 copyType = ihipCommandCopyH2D;
             } else if (kind == hipMemcpyDeviceToHost) {
                 copyType = ihipCommandCopyD2H;
@@ -2480,25 +2632,30 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
                 copyType = ihipCommandCopyD2H;
             }
 
-            hsa_signal_t depSignal;
-            int depSignalCnt = stream->copyCommand(ihip_signal, &depSignal, copyType);
+			if(trueAsync == true){
 
-            tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle);
+                hsa_signal_t depSignal;
+                int depSignalCnt = stream->copyCommand(ihip_signal, &depSignal, copyType);
 
-            hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, device->_hsa_agent, src, device->_hsa_agent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
+                tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle);
+
+                hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, device->_hsa_agent, src, device->_hsa_agent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
 
 
-            if (hsa_status == HSA_STATUS_SUCCESS) {
-                // TODO-stream - fix release-signal calls here.
-                if (HIP_LAUNCH_BLOCKING) {
-                    tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
-                    stream->wait();
+                if (hsa_status == HSA_STATUS_SUCCESS) {
+                    // TODO-stream - fix release-signal calls here.
+                    if (HIP_LAUNCH_BLOCKING) {
+                        tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
+                        stream->wait();
+                    }
+                } else {
+                    // This path can be hit if src or dst point to unpinned host memory.
+                    // TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
+                    e = hipErrorInvalidValue;
                 }
-            } else {
-                // This path can be hit if src or dst point to unpinned host memory.
-                // TODO-stream - does async-copy fall back to sync if input pointers are not pinned?
-                e = hipErrorInvalidValue;
-            }
+			} else {
+				ihipSyncCopy(stream, dst, src, sizeBytes, kind);
+			}
         }
     } else {
         e = hipErrorInvalidValue;

tests/src/CMakeLists.txt

@@ -131,8 +131,12 @@ make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrin
 make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) 
 make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp)
 make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) 
-
+make_hip_executable (hipHostAlloc hipHostAlloc.cpp)
 make_hip_executable (hipStreamL5 hipStreamL5.cpp)
+make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
+make_hip_executable (hipHostRegister hipHostRegister.cpp)
+make_hip_executable (hipRandomMemcpyAsync hipRandomMemcpyAsync.cpp)
+
 target_link_libraries(hipMathFunctionsHost m)
 
 make_test(hip_ballot " " )
@@ -151,15 +155,15 @@ make_test(hipEnvVarDriver " " )
 make_test(hipPointerAttrib " " )
 make_test(hipMultiThreadStreams1 " " )
 make_test(hipMultiThreadStreams2 " " )
-
 make_test(hipMemcpy_simple  " " )
 make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 )
 make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 )
 make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 )
+make_test(hipHostAlloc " ")
 make_test(hipMemcpyAsync  " " )
-
+make_test(hipHostGetFlags " ")
 make_test(hipHcc  " " )
-
+make_test(hipHostRegister " ")
 make_test(hipStreamL5 " ")
-
+make_test(hipRandomMemcpyAsync " ")
 make_hipify_test(specialFunc.cu )

tests/src/hipEnvVar.cpp

@@ -108,7 +108,7 @@ int main(int argc, char **argv)
         hipSetDevice(device);
         hipDeviceProp_t devProp;
 
-        hipDeviceGetProperties(&devProp, device);
+        hipGetDeviceProperties(&devProp, device);
         if (devProp.major < 1) {
             printf("%d does not support HIP\n", device);
             return -1;

tests/src/hipGetDeviceAttribute.cpp

@@ -52,7 +52,7 @@ int main(int argc, char *argv[])
     int deviceId;
     CHECK (hipGetDevice(&deviceId));
     hipDeviceProp_t props;
-    CHECK(hipDeviceGetProperties(&props, deviceId));
+    CHECK(hipGetDeviceProperties(&props, deviceId));
     printf ("info: running on device #%d %s\n", deviceId, props.name);
 
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxThreadsPerBlock, props.maxThreadsPerBlock));

tests/src/hipHcc.cpp

@@ -39,7 +39,7 @@ int main(int argc, char *argv[])
     int deviceId;
     CHECK (hipGetDevice(&deviceId));
     hipDeviceProp_t props;
-    CHECK(hipDeviceGetProperties(&props, deviceId));
+    CHECK(hipGetDeviceProperties(&props, deviceId));
     printf ("info: running on device #%d %s\n", deviceId, props.name);
 
 #ifdef __HCC__

tests/src/hipHostAlloc.cpp

@@ -0,0 +1,64 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#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(hipGetDeviceProperties(&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();
+
+}

tests/src/hipHostGetFlags.cpp

@@ -0,0 +1,81 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include"test_common.h"
+#include<malloc.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, *D;
+float *Ad, *Bd, *Cd, *Dd;
+unsigned int FlagA, FlagB, FlagC;
+FlagA = hipHostAllocWriteCombined | hipHostAllocMapped;
+FlagB = hipHostAllocWriteCombined | hipHostAllocMapped;
+FlagC = hipHostAllocMapped;
+hipDeviceProp_t prop;
+int device;
+HIPCHECK(hipGetDevice(&device));
+HIPCHECK(hipGetDeviceProperties(&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(hipHostAlloc((void**)&D, SIZE, hipHostAllocDefault));
+
+unsigned int flagA, flagB, flagC;
+
+HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Bd, B, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Dd, D, 0));
+HIPCHECK(hipHostGetFlags(&flagA, A));
+HIPCHECK(hipHostGetFlags(&flagB, B));
+HIPCHECK(hipHostGetFlags(&flagC, C));
+
+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);
+
+HIPCHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
+HIPASSERT(C[10] == 3.0f);
+HIPASSERT(flagA == FlagA);
+HIPASSERT(flagB == FlagB);
+HIPASSERT(flagC == FlagC);
+passed();
+
+}

tests/src/hipHostRegister.cpp

@@ -0,0 +1,54 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include"test_common.h"
+#include<malloc.h>
+
+__global__ void Inc(hipLaunchParm lp, float *Ad){
+int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+Ad[tx] = Ad[tx] + float(1);
+}
+
+int main(){
+	float *A, *Ad;
+	const size_t size = N * sizeof(float);
+#ifdef __HIP_PLATFORM_NVCC__
+	A = (float*)malloc(size*2);
+#else
+	A = (float*)memalign(64, size);
+#endif
+	HIPCHECK(hipHostRegister(A, size, 0));
+
+	for(int i=0;i<N;i++){
+		A[i] = float(1);
+	}
+
+	HIPCHECK(hipMalloc(&Ad, size));
+
+	HIPCHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
+
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, 0, Ad);
+	HIPCHECK(hipDeviceSynchronize());
+
+	HIPCHECK(hipMemcpy(A, Ad, size, hipMemcpyDeviceToHost));
+
+	HIPASSERT(A[10] == 2.0f);
+	HIPCHECK(hipHostUnregister(A));
+	passed();
+}

tests/src/hipKernel.cpp

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #include"test_common.h"
 
 __global__ void Empty(hipLaunchParm lp, int param){}

tests/src/hipMemcpyAsync.cpp

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 // Test under-development.  Calls async mem-copy API, experiment with functionality.
 
 #include "hip_runtime.h"

tests/src/hipRandomMemcpyAsync.cpp

@@ -0,0 +1,99 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "hip_runtime.h"
+#include "test_common.h"
+
+#define WIDTH     1024
+#define HEIGHT    1024
+
+#define NUM       (WIDTH*HEIGHT)
+
+#define THREADS_PER_BLOCK_X  16
+#define THREADS_PER_BLOCK_Y  16
+#define THREADS_PER_BLOCK_Z  1
+
+int main() {
+
+  int *hostA;
+  int *hostB;
+
+  int *deviceA;
+  int *deviceB;
+
+  int i;
+  int errors;
+
+  hostA = (int *)malloc(NUM * sizeof(int));
+  hostB = (int *)malloc(NUM * sizeof(int));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = i;
+  }
+
+  HIPCHECK(hipMalloc((void**)&deviceA, NUM * sizeof(int)));
+  HIPCHECK(hipMalloc((void**)&deviceB, NUM * sizeof(int)));
+
+  hipStream_t s;
+  HIPCHECK(hipStreamCreate(&s));
+
+
+  // hostB -> deviceB -> hostA
+#define ASYNC 1
+#if ASYNC
+  HIPCHECK(hipMemcpyAsync(deviceB, hostB, NUM*sizeof(int), hipMemcpyHostToDevice, s));
+  HIPCHECK(hipMemcpyAsync(hostA, deviceB, NUM*sizeof(int), hipMemcpyDeviceToHost, s));
+#else
+  HIPCHECK(hipMemcpy(deviceB, hostB, NUM*sizeof(int), hipMemcpyHostToDevice));
+  HIPCHECK(hipMemcpy(hostA, deviceB, NUM*sizeof(int), hipMemcpyDeviceToHost));
+#endif
+
+  HIPCHECK(hipStreamSynchronize(s));
+  HIPCHECK(hipDeviceSynchronize());
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != (hostB[i])) {
+      errors++;
+    }
+  }
+
+  HIPCHECK(hipStreamDestroy(s));
+
+  HIPCHECK(hipFree(deviceA));
+  HIPCHECK(hipFree(deviceB));
+
+  free(hostA);
+  free(hostB);
+
+  //hipResetDefaultAccelerator();
+
+  if(errors != 0){
+    HIPASSERT(1 == 2);
+  }else{
+    passed();
+  }
+
+  return errors;
+}

tests/src/hipSimpleAtomicsTest.cpp

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 // includes, system
 #include <stdlib.h>
 #include <stdio.h>
@@ -258,7 +277,7 @@ void runTest(int argc, char **argv)
     deviceProp.minor = 0;
     int dev = 0;
 
-    hipDeviceGetProperties(&deviceProp, dev);
+    hipGetDeviceProperties(&deviceProp, dev);
 
     // Statistics about the GPU device
     printf("> GPU device has %d Multi-Processors, "

tests/src/hipStream.h

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #ifndef HIPSTREAM_H
 #define HIPSTREAM_H
 #include<hip_runtime.h>

tests/src/hipStreamL5.cpp

@@ -58,7 +58,6 @@ void test12345(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -71,7 +70,7 @@ void test12345(){
 
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Ad, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Ch, Bd, size, stream);
 	HIPCHECK(hipDeviceSynchronize());
@@ -85,7 +84,6 @@ void test13452(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -107,7 +105,7 @@ void test13452(){
 	H2D(Ad, Dh, size);
 
 	H2HAsync(Bh, Ah, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Ch, Bd, size, stream);
 	H2DAsync(Cd, Ch, size, stream);
@@ -125,7 +123,6 @@ void test14523(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const int N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -150,7 +147,7 @@ void test14523(){
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Ch, Bd, size, stream);
 	H2DAsync(Cd, Ch, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
 
 	HIPCHECK(hipDeviceSynchronize());
 
@@ -165,7 +162,6 @@ void test15234(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -189,7 +185,7 @@ void test15234(){
 	H2HAsync(Bh, Ah, size, stream);
 	D2HAsync(Ch, Ad, size, stream);
 	H2DAsync(Bd, Ch, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2DAsync(Cd, Bd, size, stream);
 
 	D2H(Eh, Cd, size);
@@ -203,7 +199,6 @@ template<typename T>
 void test23451(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -218,12 +213,12 @@ void test23451(){
 	setArray(Ah, N, T(1));
 
 	H2DAsync(Ad, Ah, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Bh, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
 	HIPCHECK(hipDeviceSynchronize());
-	//HIPASSERT(Ah[10] == Ch[10]);
+	HIPASSERT(Ah[10] + T(1) == Ch[10]);
 }
 
 template<typename T>
@@ -231,7 +226,6 @@ void test24513(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -256,7 +250,7 @@ void test24513(){
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Bh, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
 	HIPCHECK(hipDeviceSynchronize());
 
 	D2H(Eh, Cd, size);
@@ -270,7 +264,6 @@ void test25134(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch;
@@ -294,7 +287,7 @@ void test25134(){
 	H2DAsync(Ad, Ah, size, stream);
 	D2HAsync(Bh, Ad, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2DAsync(Cd, Bd, size, stream);
 
 	D2H(Eh, Cd, size);
@@ -310,7 +303,6 @@ void test21345(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh, *Ch, *Dh;
@@ -328,7 +320,7 @@ void test21345(){
 
 	H2DAsync(Ad, Ah, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Dh, Bd, size, stream);
 
@@ -343,7 +335,6 @@ void test34512(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Bh, *Ch, *Dh;
@@ -363,7 +354,7 @@ void test34512(){
 
 	H2D(Ad, Ah, size);
 
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Bh, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
@@ -380,7 +371,6 @@ void test35124(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh;
@@ -399,7 +389,7 @@ void test35124(){
 
 	H2D(Ad, Dh, size);
 
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	D2HAsync(Ah, Ad, size, stream);
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Bd, Bh, size, stream);
@@ -417,7 +407,6 @@ void test31245(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 	T *Ah, *Bh, *Ch;
 	T *Dh, *Eh;
@@ -437,7 +426,7 @@ void test31245(){
 
 	H2D(Ad, Dh, size);
 
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Bd, Bh, size, stream);
 	D2DAsync(Cd, Bd, size, stream);
@@ -457,7 +446,6 @@ void test32451(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh, *Ch;
@@ -477,7 +465,7 @@ void test32451(){
 	setArray(Eh, N, T(2));
 
 	H2D(Ad, Eh, size);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Ad);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
 	H2DAsync(Bd, Ah, size, stream);
 	D2DAsync(Cd, Bd, size, stream);
 	D2HAsync(Bh, Cd, size, stream);
@@ -494,7 +482,6 @@ template<typename T>
 void test45123(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh;
@@ -517,7 +504,7 @@ void test45123(){
 	D2HAsync(Ah, Bd, size, stream);
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Cd, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
 	D2H(Ch, Cd, size);
 	HIPCHECK(hipDeviceSynchronize());
 
@@ -529,7 +516,6 @@ template<typename T>
 void test41235(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh;
@@ -550,7 +536,7 @@ void test41235(){
 	D2DAsync(Bd, Ad, size, stream);
 	D2HAsync(Ah, Bd, size, stream);
 	H2DAsync(Cd, Ah, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
 	D2HAsync(Bh, Cd, size, stream);
 
 	HIPCHECK(hipDeviceSynchronize());
@@ -563,7 +549,6 @@ void test42351(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh, *Ch;
@@ -586,7 +571,7 @@ void test42351(){
 
 	D2DAsync(Bd, Ad, size, stream);
 	H2DAsync(Cd, Ah, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Cd);
 	D2HAsync(Bh, Cd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
 
@@ -602,7 +587,6 @@ void test43512(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh;
@@ -622,7 +606,7 @@ void test43512(){
 	H2D(Ad, Dh, size);
 
 	D2DAsync(Bd, Ad, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2HAsync(Ah, Bd, size, stream);
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Cd, Bh, size, stream);
@@ -637,7 +621,6 @@ void test51234(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh;
@@ -659,7 +642,7 @@ void test51234(){
 	D2HAsync(Ah, Ad, size, stream);
 	H2HAsync(Bh, Ah, size, stream);
 	H2DAsync(Bd, Bh, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2DAsync(Cd, Bd, size, stream);
 
 	D2H(Ch, Cd, size);
@@ -673,7 +656,6 @@ template<typename T>
 void test52341(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh, *Ch;
@@ -696,7 +678,7 @@ void test52341(){
 
 	D2HAsync(Ah, Ad, size, stream);
 	H2DAsync(Bd, Ah, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2DAsync(Cd, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
 
@@ -712,7 +694,6 @@ template<typename T>
 void test53412(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
-	const size_t N = 1000;
 	const size_t size = sizeof(T) * N;
 
 	T *Ah, *Bh, *Ch, *Dh;
@@ -739,7 +720,7 @@ void test53412(){
 	H2D(Bd, Eh, size);
 
 	D2HAsync(Ah, Ad, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Bd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Bd);
 	D2DAsync(Cd, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
 	H2DAsync(Dd, Ch, size, stream);
@@ -757,7 +738,6 @@ void test54123(){
 	hipStream_t stream;
 	HIPCHECK(hipStreamCreate(&stream));
 
-	const size_t N = 1000;
 	const size_t size = N * sizeof(T);
 
 	T *Ah, *Bh, *Ch;
@@ -787,7 +767,7 @@ void test54123(){
 	D2DAsync(Cd, Bd, size, stream);
 	H2HAsync(Ch, Bh, size, stream);
 	H2DAsync(Dd, Ch, size, stream);
-	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Dd);
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Dd);
 
 	D2H(Fh, Cd, size);
 	D2H(Gh, Dd, size);

tests/src/hip_anyall.cpp

@@ -41,7 +41,7 @@ __global__ void
 int main(int argc, char *argv[])
 { int warpSize, pshift;
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   if(strncmp(devProp.name,"Fiji",1)==0) 
 { warpSize =64;
   pshift =6;

tests/src/hip_ballot.cpp

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #include <iostream>
 
 #include <hip_runtime.h>
@@ -21,7 +40,7 @@ __global__ void
 int main(int argc, char *argv[])
 { int warpSize, pshift;
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   
   if(strncmp(devProp.name,"Fiji",1)==0)  
   {warpSize = 64; pshift =6;}

tests/src/hip_brev.cpp

@@ -94,7 +94,7 @@ int main() {
   unsigned long long int* deviceD;
 
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   cout << " System minor " << devProp.minor << endl;
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;

tests/src/hip_clz.cpp

@@ -118,7 +118,7 @@ int main() {
   long long int* deviceH;
 
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   cout << " System minor " << devProp.minor << endl;
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;

tests/src/hip_ffs.cpp

@@ -89,7 +89,7 @@ int main() {
   unsigned long long int* deviceD;
   
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   cout << " System minor " << devProp.minor << endl;
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;

tests/src/hip_popc.cpp

@@ -86,7 +86,7 @@ int main() {
   unsigned long long int* deviceD;
 
   hipDeviceProp_t devProp;
-  hipDeviceGetProperties(&devProp, 0);
+  hipGetDeviceProperties(&devProp, 0);
   cout << " System minor " << devProp.minor << endl;
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;

tests/src/specialFunc.cu

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 //Test to ensure hipify runs correctly.
 // Hipify may report warnings for some missing/unsupported functions
 

tests/src/test_common.cpp

@@ -144,7 +144,7 @@ unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N)
     int device;
     HIPCHECK(hipGetDevice(&device));
     hipDeviceProp_t props;
-    HIPCHECK(hipDeviceGetProperties(&props, device));
+    HIPCHECK(hipGetDeviceProperties(&props, device));
 
     unsigned blocks = props.multiProcessorCount * blocksPerCU;
     if (blocks * threadsPerBlock > N) {

tests/src/test_common.h

@@ -1,3 +1,22 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
 #include <iostream>
 #include <sys/time.h>
 #include <stddef.h>

util/gedit/hip.lang

@@ -226,8 +226,8 @@
           <keyword>hipD3D9SetDirect3DDevice</keyword>
           <keyword>hipD3D9UnmapResources</keyword>
           <keyword>hipD3D9UnregisterResource</keyword>
-          <keyword>hipDeviceGetProperties</keyword>
-					<keyword>hipDeviceSynchronize</keyword>
+          <keyword>hipGetDeviceProperties</keyword>
+		  <keyword>hipDeviceSynchronize</keyword>
           <keyword>hipEventCreate</keyword>
           <keyword>hipEventDestroy</keyword>
           <keyword>hipEventElapsedTime</keyword>

util/vim/hip.vim

@@ -92,7 +92,7 @@ syn keyword hipFunctionName hipD3D9ResourceSetMapFlags
 syn keyword hipFunctionName hipD3D9SetDirect3DDevice 
 syn keyword hipFunctionName hipD3D9UnmapResources 
 syn keyword hipFunctionName hipD3D9UnregisterResource 
-syn keyword hipFunctionName hipDeviceGetProperties 
+syn keyword hipFunctionName hipGetDeviceProperties 
 syn keyword hipFunctionName hipDeviceSynchronize 
 syn keyword hipFunctionName hipDeviceReset 
 syn keyword hipFunctionName hipEventCreate