Merge 'master' into 'amd-master'

Change-Id: I7782abb64120ea7da07d0cd3af4dbf1d60ba4f09

bin/hipcc

@@ -106,6 +106,7 @@ if ($HIP_PLATFORM eq "hcc") {
 
     #### GCC system includes workaround ####
     $HCC_WA_FLAGS = " ";
+    $HOST_OSNAME= `cat /etc/os-release | grep "^ID\=" | cut -d= -f2 | tr -d '\n'`;
     if ($HCC_VERSION_MAJOR eq 1) {
         my $GCC_CUR_VER = `gcc -dumpversion`;
         my $GPP_CUR_VER = `g++ -dumpversion`;
@@ -116,7 +117,8 @@ if ($HIP_PLATFORM eq "hcc") {
 
         # Only include the libstdc++ headers and libraries flags explicitly if the g++ is older than version 5.
         # That's because HCC already uses libstdc++ by default if a newer g++/libstdc++ is available
-        if (${GCC_CUR_VER} eq ${GPP_CUR_VER} and $GPP_VER_FIELDS[0] < 5) {
+        # Cent OS 7 and RHEL 7.4 cannot use libstdc++ for compilation, default to libc++
+        if (${GCC_CUR_VER} eq ${GPP_CUR_VER} and $GPP_VER_FIELDS[0] < 5 and ($HOST_OSNAME ne "\"centos\"") and ($HOST_OSNAME ne "\"rhel\"")) {
             $HCC_WA_FLAGS .= " -stdlib=libstdc++  -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
             # Add C++ libs for GCC.
             $HIPLDFLAGS .= " -lstdc++";
@@ -124,7 +126,6 @@ if ($HIP_PLATFORM eq "hcc") {
     }
 
     # Force -stdlib=libc++ on UB14.04
-    $HOST_OSNAME= `cat /etc/os-release | grep "^ID\=" | cut -d= -f2 | tr -d '\n'`;
     $HOST_OSVER= `cat /etc/os-release | grep "^VERSION_ID\=" | cut -d= -f2 | tr -d '\n'`;
     if ($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"14.04\"") {
         $HIPCXXFLAGS .= " -stdlib=libc++";

bin/hipify-perl

@@ -247,6 +247,7 @@ while (@ARGV) {
         $ft{'err'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g;
         $ft{'err'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g;
         $ft{'err'} += s/\bcudaErrorInvalidDevice\b/hipErrorInvalidDevice/g;
+        $ft{'err'} += s/\bcudaErrorInvalidDevicePointer\b/hipErrorInvalidDevicePointer/g;
         $ft{'err'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g;
         $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g;
         $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
@@ -280,6 +281,7 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g;
 
         $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g;
+        $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttribute_t/g;
 
         $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g;
         $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
@@ -339,6 +341,7 @@ while (@ARGV) {
         $ft{'event'} += s/\bcudaEventElapsedTime\b/hipEventElapsedTime/g;
         $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g;
         $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g;
+        $ft{'event'} += s/\bcudaEventQuery\b/hipEventQuery/g;
 
         #--------
         # Streams
@@ -489,15 +492,15 @@ while (@ARGV) {
             my $kernelName;
 
             # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, $6, /g;
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g;
             $kernelName = $1 if $k;
 
             # Handle the <<numBlocks, blockDim, sharedSize>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, 0, /g;
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g;
             $kernelName = $1 if $k;
 
             # Handle the <<numBlocks, blockDim>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), 0, 0, /g;
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g;
             $kernelName = $1 if $k;
 
             $ft{'kern'} += $k;

docs/markdown/hip_kernel_language.md

@@ -699,8 +699,18 @@ for (int i=0; i<16; i++) ...
 
 ## In-Line Assembly
 
-In-line assembly, including in-line PTX, in-line HSAIL and in-line GCN ISA, is not supported. Users who need these features should employ conditional compilation to provide different functionally equivalent implementations on each target platform.
+GCN ISA In-line assembly, is supported. For example:
 
+```
+asm volatile ("v_mac_f32_e32 %0, %2, %3" : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i]));
+```
+
+We insert the GCN isa into the kernel using `asm()` Assembler statement.
+`volatile` keyword is used so that the optimizers must not change the number of volatile operations or change their order of execution relative to other volatile operations.
+`v_mac_f32_e32` is the GCN instruction, for more information please refer - [AMD GCN3 ISA architecture manual](http://gpuopen.com/compute-product/amd-gcn3-isa-architecture-manual/)
+Index for the respective operand in the ordered fashion is provided by `%` followed by position in the list of operands
+`"v"` is the constraint code (for target-specific AMDGPU) for 32-bit VGPR register, for more info please refer - [Supported Constraint Code List for AMDGPU](https://llvm.org/docs/LangRef.html#supported-constraint-code-list)
+Output Constraints are specified by an `"="` prefix as shown above ("=v"). This indicate that assemby will write to this operand, and the operand will then be made available as a return value of the asm expression. Input constraints do not have a prefix - just the constraint code. The constraint string of `"0"` says to use the assigned register for output as an input as well (it being the 0'th constraint).
 
 ## C++ Support
 The following C++ features are not supported:

include/hip/hcc_detail/hip_runtime.h

@@ -99,7 +99,7 @@ extern int HIP_TRACE_API;
 // TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call.
 #if __HIP_DEVICE_COMPILE__ == 1
    #undef assert
-   #define assert(COND) { if (COND) {} }
+   #define assert(COND) { if (!COND) {abort();} }
 #endif
 
 

include/hip/hip_runtime_api.h

@@ -161,6 +161,7 @@ typedef enum hipError_t {
     hipErrorProfilerNotInitialized  = 6,
     hipErrorProfilerAlreadyStarted  = 7,
     hipErrorProfilerAlreadyStopped  = 8,
+    hipErrorInsufficientDriver      = 35,
     hipErrorInvalidImage            = 200,
     hipErrorInvalidContext          = 201,  ///< Produced when input context is invalid.
     hipErrorContextAlreadyCurrent   = 202,

include/hip/nvcc_detail/hip_runtime.h

@@ -111,6 +111,8 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(__VA_ARGS__);\
 
 #ifdef __HIP_DEVICE_COMPILE__
 #define abort() {asm("trap;");}
+#undef assert
+#define assert(COND) { if (!COND) {abort();} }
 #endif
 
 #endif

include/hip/nvcc_detail/hip_runtime_api.h

@@ -31,6 +31,13 @@ THE SOFTWARE.
 extern "C" {
 #endif
 
+#ifdef __cplusplus
+  #define __dparm(x) \
+          = x
+#else
+  #define __dparm(x)
+#endif
+ 
     //TODO -move to include/hip_runtime_api.h as a common implementation.
 /**
 * Memory copy types
@@ -102,20 +109,20 @@ typedef cudaEvent_t hipEvent_t;
 typedef cudaStream_t hipStream_t;
 typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
 typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
-typedef cudaLimit hipLimit_t;
-typedef cudaFuncCache hipFuncCache_t;
+typedef enum cudaLimit hipLimit_t;
+typedef enum cudaFuncCache hipFuncCache_t;
 typedef CUcontext hipCtx_t;
-typedef CUsharedconfig hipSharedMemConfig;
+typedef cudaSharedMemConfig hipSharedMemConfig;
 typedef CUfunc_cache hipFuncCache;
 typedef CUjit_option hipJitOption;
 typedef CUdevice hipDevice_t;
 typedef CUmodule hipModule_t;
 typedef CUfunction hipFunction_t;
 typedef CUdeviceptr hipDeviceptr_t;
-typedef cudaChannelFormatKind hipChannelFormatKind;
-typedef cudaChannelFormatDesc hipChannelFormatDesc;
-typedef cudaTextureReadMode hipTextureReadMode;
-typedef cudaArray hipArray;
+typedef enum cudaChannelFormatKind hipChannelFormatKind;
+typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
+typedef enum cudaTextureReadMode hipTextureReadMode;
+typedef struct cudaArray hipArray;
 
 // Flags that can be used with hipStreamCreateWithFlags
 #define hipStreamDefault            cudaStreamDefault
@@ -124,6 +131,11 @@ typedef cudaArray hipArray;
 //typedef cudaChannelFormatDesc hipChannelFormatDesc;
 #define hipChannelFormatDesc cudaChannelFormatDesc
 
+//adding code for hipmemSharedConfig
+#define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault
+#define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
+#define hipSharedMemBankSizeEightByte cudaSharedMemBankSizeEightByte
+
 inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
 switch(cuError) {
     case cudaSuccess                             : return hipSuccess;
@@ -187,7 +199,7 @@ switch(hError) {
 }
 }
 
-inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
+inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
     switch(kind) {
     case hipMemcpyHostToHost:
         return cudaMemcpyHostToHost;
@@ -250,7 +262,7 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla
 	return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
 }
 
-inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) {
+inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) {
   return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
 }
 
@@ -289,8 +301,8 @@ inline static hipError_t hipSetDevice(int device) {
 
 inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
 {
-    cudaDeviceProp cdprop;
-    memset(&cdprop,0x0,sizeof(cudaDeviceProp));
+    struct cudaDeviceProp cdprop;
+    memset(&cdprop,0x0,sizeof(struct cudaDeviceProp));
     cdprop.major= prop->major;
     cdprop.minor = prop->minor;
     cdprop.totalGlobalMem = prop->totalGlobalMem ;
@@ -351,25 +363,24 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
 }
 
 
-inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream=0) {
+inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {
   return hipCUDAErrorTohipError(cudaMemcpyAsync(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind), stream));
 }
 
-
-inline static hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeBytes, size_t offset = 0, hipMemcpyKind copyType = hipMemcpyHostToDevice) {
+inline static hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind copyType __dparm(hipMemcpyHostToDevice)) {
 	return hipCUDAErrorTohipError(cudaMemcpyToSymbol(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType)));
 }
 
-inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind copyType, hipStream_t stream = 0) {
+inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind copyType, hipStream_t stream __dparm(0)) {
     return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType), stream));
 }
 
-inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset = 0, hipMemcpyKind kind = hipMemcpyDeviceToHost)
+inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost))
 {
     return hipCUDAErrorTohipError(cudaMemcpyFromSymbol(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind)));
 }
 
-inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream = 0)
+inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0))
 {
     return hipCUDAErrorTohipError(cudaMemcpyFromSymbolAsync(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream));
 }
@@ -438,7 +449,7 @@ inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
     return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
 }
 
-inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream = 0) {
+inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream __dparm(0)) {
     return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream));
 }
 
@@ -449,7 +460,7 @@ inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char  value,
 
 inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device)
 {
-	cudaDeviceProp cdprop;
+	struct cudaDeviceProp cdprop;
 	cudaError_t cerror;
 	cerror = cudaGetDeviceProperties(&cdprop,device);
 	strncpy(p_prop->name,cdprop.name, 256);
@@ -510,7 +521,7 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int dev
 
 inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
 {
-    cudaDeviceAttr cdattr;
+    enum cudaDeviceAttr cdattr;
     cudaError_t cerror;
 
     switch (attr) {
@@ -586,7 +597,7 @@ inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
 }
 
 inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr){
-	cudaPointerAttributes cPA;
+	struct cudaPointerAttributes cPA;
 	hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr));
 	if(err == hipSuccess){
 		switch (cPA.memoryType){
@@ -617,7 +628,7 @@ inline static hipError_t hipEventCreate( hipEvent_t* event)
     return hipCUDAErrorTohipError(cudaEventCreate(event));
 }
 
-inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream = NULL)
+inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream __dparm(NULL))
 {
     return hipCUDAErrorTohipError(cudaEventRecord(event,stream));
 }
@@ -750,18 +761,18 @@ inline static hipError_t hipMemcpyPeer ( void* dst, int  dstDevice, const void*
     return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count));
 }
 
-inline static hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count, hipStream_t stream=0 )
+inline static hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count, hipStream_t stream __dparm(0))
 {
     return hipCUDAErrorTohipError(cudaMemcpyPeerAsync(dst, dstDevice, src, srcDevice, count, stream));
 }
 
 // Profile APIs:
-inline hipError_t hipProfilerStart()
+inline static hipError_t hipProfilerStart()
 {
     return hipCUDAErrorTohipError(cudaProfilerStart());
 }
 
-inline hipError_t hipProfilerStop()
+inline static hipError_t hipProfilerStop()
 {
     return hipCUDAErrorTohipError(cudaProfilerStop());
 }
@@ -833,12 +844,12 @@ inline static hipError_t  hipCtxSetCacheConfig (hipFuncCache cacheConfig)
 
 inline static hipError_t  hipCtxSetSharedMemConfig (hipSharedMemConfig config)
 {
-    return hipCUResultTohipError(cuCtxSetSharedMemConfig(config));
+    return hipCUResultTohipError(cuCtxSetSharedMemConfig((CUsharedconfig)config));
 }
 
 inline static hipError_t  hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
 {
-    return hipCUResultTohipError(cuCtxGetSharedMemConfig(pConfig));
+    return hipCUResultTohipError(cuCtxGetSharedMemConfig((CUsharedconfig *)pConfig));
 }
 
 inline static hipError_t  hipCtxSynchronize ( void )
@@ -881,6 +892,16 @@ inline static hipError_t hipDeviceGetByPCIBusId(int* device, const char *pciBusI
     return hipCUDAErrorTohipError(cudaDeviceGetByPCIBusId(device, pciBusId));
 }
 
+inline static hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig *config)
+{
+    return hipCUDAErrorTohipError(cudaDeviceGetSharedMemConfig(config));
+}
+ 
+inline static hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config)
+{
+    return hipCUDAErrorTohipError(cudaDeviceSetSharedMemConfig(config));
+}
+
 inline static hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit)
 {
     return hipCUDAErrorTohipError(cudaDeviceGetLimit(pValue, limit));

samples/2_Cookbook/10_inline_asm/Readme.md

@@ -27,10 +27,23 @@ We will be using the Simple Matrix Transpose application from the our very first
 
 ## asm() Assembler statement
 
-We insert the GCN isa into the kernel using asm() Assembler statement. In the same sourcecode, we used for MatrixTranspose. We'll add the following:
+In the same sourcecode, we used for MatrixTranspose. We'll add the following:
 
 `  asm volatile ("v_mov_b32_e32 %0, %1" : "=v" (out[x*width + y]) : "v" (in[y*width + x]));                    `
 
+GCN ISA In-line assembly, is supported. For example:
+
+```
+asm volatile ("v_mac_f32_e32 %0, %2, %3" : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i]));
+```
+
+We insert the GCN isa into the kernel using `asm()` Assembler statement.
+`volatile` keyword is used so that the optimizers must not change the number of volatile operations or change their order of execution relative to other volatile operations.
+`v_mac_f32_e32` is the GCN instruction, for more information please refer - [AMD GCN3 ISA architecture manual](http://gpuopen.com/compute-product/amd-gcn3-isa-architecture-manual/)
+Index for the respective operand in the ordered fashion is provided by `%` followed by position in the list of operands
+`"v"` is the constraint code (for target-specific AMDGPU) for 32-bit VGPR register, for more info please refer - [Supported Constraint Code List for AMDGPU](https://llvm.org/docs/LangRef.html#supported-constraint-code-list)
+Output Constraints are specified by an `"="` prefix as shown above ("=v"). This indicate that assemby will write to this operand, and the operand will then be made available as a return value of the asm expression. Input constraints do not have a prefix - just the constraint code. The constraint string of `"0"` says to use the assigned register for output as an input as well (it being the 0'th constraint).
+
 ## How to build and run:
 Use the make command and execute it using ./exe
 Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.

src/hip_hcc.cpp

@@ -89,6 +89,12 @@ int HIP_HOST_COHERENT = 1;
 int HIP_SYNC_HOST_ALLOC = 1;
 
 
+int HIP_INIT_ALLOC=-1;
+int HIP_SYNC_STREAM_WAIT = 0;
+int HIP_FORCE_NULL_STREAM=0;
+
+
+
 #if (__hcc_workweek__ >= 17300)
 // Make sure we have required bug fix in HCC
 // Perform resolution on the GPU:
@@ -267,8 +273,6 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int
         case hipDeviceScheduleBlockingSync  : _scheduleMode = Yield; break;
         default:_scheduleMode = Auto;
     };
-
-
 };
 
 
@@ -1248,7 +1252,11 @@ void HipReadEnv()
     READ_ENV_I(release, HIP_FAIL_SOC, 0, "Fault on Sub-Optimal-Copy, rather than use a slower but functional implementation.  Bit 0x1=Fail on async copy with unpinned memory.  Bit 0x2=Fail peer copy rather than use staging buffer copy");
 
     READ_ENV_I(release, HIP_SYNC_HOST_ALLOC, 0, "Sync before and after all host memory allocations.  May help stability");
+    READ_ENV_I(release, HIP_INIT_ALLOC, 0, "If not -1, initialize allocated memory to specified byte");
     READ_ENV_I(release, HIP_SYNC_NULL_STREAM, 0, "Synchronize on host for null stream submissions");
+    READ_ENV_I(release, HIP_FORCE_NULL_STREAM,  0, "Force all stream allocations to secretly return the null stream");
+
+    READ_ENV_I(release, HIP_SYNC_STREAM_WAIT, 0, "hipStreamWaitEvent will synchronize to host");
 
 
     READ_ENV_I(release, HIP_HOST_COHERENT, 0, "If set, all host memory will be allocated as fine-grained system memory.  This allows threadfence_system to work but prevents host memory from being cached on GPU which may have performance impact.");

src/hip_hcc_internal.h

@@ -65,8 +65,12 @@ extern int HIP_HIDDEN_FREE_MEM;
 //---
 // Chicken bits for disabling functionality to work around potential issues:
 extern int HIP_SYNC_HOST_ALLOC;
+extern int HIP_SYNC_STREAM_WAIT;
 
 extern int HIP_SYNC_NULL_STREAM;
+extern int HIP_INIT_ALLOC;
+extern int HIP_FORCE_NULL_STREAM;
+
 
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;

src/hip_memory.cpp

@@ -119,6 +119,11 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool
     tprintf(DB_MEM, " alloc %s ptr:%p-%p size:%zu on dev:%d\n",
             msg, ptr, static_cast<char*>(ptr)+sizeBytes, sizeBytes, device->_deviceId);
 
+    if (HIP_INIT_ALLOC != -1) {
+        // TODO , dont' call HIP API directly here:
+        hipMemset(ptr, HIP_INIT_ALLOC, sizeBytes);
+    }	
+
     if (ptr != nullptr) {
         int r = sharePtr(ptr, ctx, shareWithAll, hipFlags);
         if (r != 0) {
@@ -407,83 +412,89 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
     HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
     HIP_SET_DEVICE();
     hipError_t  hip_status = hipSuccess;
+    if(width > 0) {
+        auto ctx = ihipGetTlsDefaultCtx();
 
-    auto ctx = ihipGetTlsDefaultCtx();
+        *array = (hipArray*)malloc(sizeof(hipArray));
+        array[0]->type  = flags;
+        array[0]->width = width;
+        array[0]->height = height;
+        array[0]->depth = 1;
+        array[0]->desc = *desc;
 
-    *array = (hipArray*)malloc(sizeof(hipArray));
-    array[0]->type  = flags;
-    array[0]->width = width;
-    array[0]->height = height;
-    array[0]->depth = 1;
-    array[0]->desc = *desc;
+        void ** ptr = &array[0]->data;
 
-    void ** ptr = &array[0]->data;
+        if (ctx) {
+            const unsigned am_flags = 0;
+            size_t size = width;
+            if(height > 0) {
+                size = size * height;
+            }
 
-    if (ctx) {
-        const unsigned am_flags = 0;
-        const size_t size = width*height;
+            size_t allocSize = 0;
+            switch(desc->f) {
+                case hipChannelFormatKindSigned:
+                    allocSize = size * sizeof(int);
+                    break;
+                case hipChannelFormatKindUnsigned:
+                    allocSize = size * sizeof(unsigned int);
+                    break;
+                case hipChannelFormatKindFloat:
+                    allocSize = size * sizeof(float);
+                    break;
+                case hipChannelFormatKindNone:
+                    allocSize = size * sizeof(size_t);
+                    break;
+                default:
+                    hip_status = hipErrorUnknown;
+                    break;
+            }
+            hc::accelerator acc = ctx->getDevice()->_acc;
+            hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
 
-        size_t allocSize = 0;
-        switch(desc->f) {
-            case hipChannelFormatKindSigned:
-                allocSize = size * sizeof(int);
-                break;
-            case hipChannelFormatKindUnsigned:
-                allocSize = size * sizeof(unsigned int);
-                break;
-            case hipChannelFormatKindFloat:
-                allocSize = size * sizeof(float);
-                break;
-            case hipChannelFormatKindNone:
-                allocSize = size * sizeof(size_t);
+            size_t allocGranularity = 0;
+            hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
+            hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
+
+            hsa_ext_image_descriptor_t imageDescriptor;
+
+            imageDescriptor.width = width;
+            imageDescriptor.height = height;
+            imageDescriptor.depth = 0;
+            imageDescriptor.array_size = 0;
+            switch (flags) {
+            case hipArrayLayered:
+            case hipArrayCubemap:
+            case hipArraySurfaceLoadStore:
+            case hipArrayTextureGather:
+                assert(0);
                 break;
+            case hipArrayDefault:
             default:
-                hip_status = hipErrorUnknown;
+                imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
                 break;
-        }
-        hc::accelerator acc = ctx->getDevice()->_acc;
-        hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+            }
+            hsa_ext_image_channel_order_t channelOrder;
+            hsa_ext_image_channel_type_t channelType;
+            getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
+            imageDescriptor.format.channel_order = channelOrder;
+            imageDescriptor.format.channel_type = channelType;
 
-        size_t allocGranularity = 0;
-        hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
-        hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
+            hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+            hsa_ext_image_data_info_t imageInfo;
+            hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
+            size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
 
-        hsa_ext_image_descriptor_t imageDescriptor;
+            *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
+            if (size && (*ptr == NULL)) {
+                hip_status = hipErrorMemoryAllocation;
+            }
 
-        imageDescriptor.width = width;
-        imageDescriptor.height = height;
-        imageDescriptor.depth = 0;
-        imageDescriptor.array_size = 0;
-        switch (flags) {
-        case hipArrayLayered:
-        case hipArrayCubemap:
-        case hipArraySurfaceLoadStore:
-        case hipArrayTextureGather:
-            assert(0);
-            break;
-        case hipArrayDefault:
-        default:
-            imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
-            break;
-        }
-        hsa_ext_image_channel_order_t channelOrder;
-        hsa_ext_image_channel_type_t channelType;
-        getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
-        imageDescriptor.format.channel_order = channelOrder;
-        imageDescriptor.format.channel_type = channelType;
-
-        hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
-        hsa_ext_image_data_info_t imageInfo;
-        hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
-        size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
-
-        *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
-        if (size && (*ptr == NULL)) {
+        } else {
             hip_status = hipErrorMemoryAllocation;
         }
-
     } else {
-        hip_status = hipErrorMemoryAllocation;
+         hip_status = hipErrorInvalidValue;
     }
 
     return ihipLogStatus(hip_status);
@@ -1180,6 +1191,7 @@ ihipMemsetKernel(hipStream_t stream,
 
 }
 
+
 // TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
 hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t stream )
 {

src/hip_stream.cpp

@@ -38,21 +38,26 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
     hipError_t e = hipSuccess;
 
     if (ctx) {
-        hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
-        // TODO - se try-catch loop to detect memory exception?
-        //
-        //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
-        //This matches CUDA stream behavior:
+        if (HIP_FORCE_NULL_STREAM) {
+            *stream = 0; 
+        } else {
+            hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
-        {
-            // Obtain mutex access to the device critical data, release by destructor
-            LockedAccessor_CtxCrit_t  ctxCrit(ctx->criticalData());
+            // TODO - se try-catch loop to detect memory exception?
+            //
+            //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
+            //This matches CUDA stream behavior:
 
-            auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
+            {
+                // Obtain mutex access to the device critical data, release by destructor
+                LockedAccessor_CtxCrit_t  ctxCrit(ctx->criticalData());
 
-            ctxCrit->addStream(istream);
-            *stream = istream;
+                auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
+
+                ctxCrit->addStream(istream);
+                *stream = istream;
+            }
         }
 
         tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
@@ -93,18 +98,15 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
     } else if (event->_state != hipEventStatusUnitialized) {
 
-        if (stream != hipStreamNull) {
-
+        if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
+            // conservative wait on host for the specified event to complete:
+            event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
+        } else {
+            stream = ihipSyncAndResolveStream(stream);
             // This will user create_blocking_marker to wait on the specified queue.
             stream->locked_streamWaitEvent(event);
-
-        } else {
-            // TODO-hcc Convert to use create_blocking_marker(...) functionality.
-            // Currently we have a super-conservative version of this - block on host, and drain the queue.
-            // This should create a barrier packet in the target queue.
-            // TODO-HIP_SYNC_NULL_STREAM
-            stream->locked_wait();
         }
+
     } // else event not recorded, return immediately and don't create marker.
 
     return ihipLogStatus(e);
@@ -122,15 +124,14 @@ hipError_t hipStreamQuery(hipStream_t stream)
         stream =  device->_defaultStream;
     }
 
-    int pendingOps = 0;
+    bool isEmpty = 0;
 
     {
         LockedAccessor_StreamCrit_t crit(stream->_criticalData);
-        pendingOps = crit->_av.get_pending_async_ops();
+        isEmpty = crit->_av.get_is_empty();
     }
 
-
-    hipError_t e = (pendingOps > 0) ? hipErrorNotReady : hipSuccess;
+    hipError_t e = isEmpty ? hipSuccess : hipErrorNotReady ;
 
     return ihipLogStatus(e);
 }
@@ -170,7 +171,9 @@ hipError_t hipStreamDestroy(hipStream_t stream)
 
     //--- Drain the stream:
     if (stream == NULL) {
-        e = hipErrorInvalidResourceHandle; // TODO - review - what happens if try to destroy null stream
+        if (!HIP_FORCE_NULL_STREAM) {
+            e = hipErrorInvalidResourceHandle; 
+        } 
     } else {
         stream->locked_wait();
 

tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp

@@ -446,9 +446,41 @@ int main(int argc, char *argv[])
 
 
     if (p_tests & 0x1000) {
-        printf ("==> Test 0x1000 try null stream\n"); 
-        hipStreamQuery(0/* try null stream*/);
+        printf ("==> Test 0x1000 simple null stream tests\n"); 
 
+        // try some null stream:
+        hipStreamQuery(0);
+
+
+        hipStream_t s1;
+        hipEvent_t e1;
+
+        {
+            // stream null waits on event in s1 stream:
+            HIPCHECK(hipStreamCreate(&s1));
+            HIPCHECK(hipEventCreate(&e1));
+
+            HIPCHECK(hipEventRecord(e1, s1))
+
+            HIPCHECK(hipStreamWaitEvent(hipStream_t(0), e1, 0/*flags*/));
+            
+            HIPCHECK(hipStreamDestroy(s1));
+            HIPCHECK(hipEventDestroy(e1));
+        }
+
+        {
+            // stream s1 waits on event in null stream:
+            HIPCHECK(hipStreamCreate(&s1));
+            HIPCHECK(hipEventCreate(&e1));
+
+            HIPCHECK(hipEventRecord(e1, hipStream_t(0)))
+
+            HIPCHECK(hipStreamWaitEvent(s1, e1, 0/*flags*/));
+            
+            HIPCHECK(hipStreamDestroy(s1));
+            HIPCHECK(hipEventDestroy(e1));
+        }
+        
     }
 
 
@@ -471,8 +503,8 @@ int main(int argc, char *argv[])
     }
 
 
-    {
-        printf ("test: alternating memcpy/count-reverse followed by event\n");
+    if (p_tests & 0x4000 ) {
+        printf ("test: %x alternating memcpy/count-reverse followed by event\n", p_tests);
         RUN_SYNC_TEST(0x4000, streamersDev0, sync_queryAllUntilComplete(streamersDev0),  true);
         RUN_SYNC_TEST(0x8000, streamersDev0, sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false),  true);
     }