diff --git a/bin/hipcc b/bin/hipcc index 624777ab25..8f18ab0d8e 100755 --- a/bin/hipcc +++ b/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++"; diff --git a/bin/hipify-perl b/bin/hipify-perl index 4b6e6e89b9..a37f728bd1 100755 --- a/bin/hipify-perl +++ b/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 <>> 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 <>> 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 <>> 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; diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index a4f4d5d47f..094d7531e8 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/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: diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 1092e22c47..2386ea08cb 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/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 diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 0cdace0e99..97733964bf 100644 --- a/include/hip/hip_runtime_api.h +++ b/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, diff --git a/include/hip/nvcc_detail/hip_runtime.h b/include/hip/nvcc_detail/hip_runtime.h index 8c08f3d151..fca9ab6e39 100644 --- a/include/hip/nvcc_detail/hip_runtime.h +++ b/include/hip/nvcc_detail/hip_runtime.h @@ -111,6 +111,8 @@ kernelName<<>>(__VA_ARGS__);\ #ifdef __HIP_DEVICE_COMPILE__ #define abort() {asm("trap;");} +#undef assert +#define assert(COND) { if (!COND) {abort();} } #endif #endif diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 502d205ac4..1103bb2274 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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)); diff --git a/samples/2_Cookbook/10_inline_asm/Readme.md b/samples/2_Cookbook/10_inline_asm/Readme.md index 0e64fe9c6e..7d0301bc74 100644 --- a/samples/2_Cookbook/10_inline_asm/Readme.md +++ b/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. diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index dc72714e3e..445e74cbeb 100644 --- a/src/hip_hcc.cpp +++ b/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."); diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 56ca37b3e2..4b7e533a4c 100644 --- a/src/hip_hcc_internal.h +++ b/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; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 4b3a4fcb12..a8324c5729 100644 --- a/src/hip_memory.cpp +++ b/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(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(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(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(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(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 ) { diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index 51aeb01412..7dd6efd39c 100644 --- a/src/hip_stream.cpp +++ b/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(); diff --git a/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 9bbd43828c..f5b1b79550 100644 --- a/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/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); }