diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 37d3534cd8..e418a184ae 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -217,13 +217,14 @@ endif() # Install .version install(FILES ${PROJECT_BINARY_DIR}/.version DESTINATION bin) -# Install src, bin, include if necessary +# Install src, bin, include & cmake if necessary execute_process(COMMAND test ${CMAKE_INSTALL_PREFIX} -ef ${CMAKE_CURRENT_SOURCE_DIR} RESULT_VARIABLE INSTALL_SOURCE) if(NOT ${INSTALL_SOURCE} EQUAL 0) install(DIRECTORY src DESTINATION .) install(DIRECTORY bin DESTINATION . USE_SOURCE_PERMISSIONS) install(DIRECTORY include DESTINATION .) + install(DIRECTORY cmake DESTINATION .) endif() ############################# diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index 060e70e519..7858ef58eb 100644 --- a/projects/hip/CONTRIBUTING.md +++ b/projects/hip/CONTRIBUTING.md @@ -94,7 +94,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl - Code Indentation: - Tabs should be expanded to spaces. - Use 4 spaces indendation. -- Capitaliziation and Naming +- Capitalization and Naming - Prefer camelCase for HIP interfaces and internal symbols. Note HCC uses _ for separator. This guideline is not yet consistently followed in HIP code - eventual compliance is aspirational. - Member variables should begin with a leading "_". This allows them to be easily distinguished from other variables or functions. @@ -110,6 +110,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl doFooElse(); } ''' + - namespace should be on same line as { and separated by a space. - Single-line if statement should still use {/} pair (even though C++ does not require). - Miscellaneous - All references in function parameter lists should be const. @@ -120,6 +121,8 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl - HIP_INIT_API() should be placed at the start of each top-level HIP API. This function will make sure the HIP runtime is initialized, and also constructs an appropriate API string for tracing and CodeXL marker tracing. The arguments to HIP_INIT_API should match those of the parent fucntion. +- ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code. The error code + is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly. #### Presubmit Testing: diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 8fde7f426d..c3d46bfcce 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -168,6 +168,8 @@ my $hasCU = 0; # options contain a cu-style file (HCC must force recogni my $needHipHcc = ($HIP_PLATFORM eq 'hcc'); # set if we need to link hip_hcc.o from src tree. (some builds, ie cmake, provide their own) my $printHipVersion = 0; # print HIP version my $runCmd = 1; +my $buildDeps = 0; + my @options = (); my @inputs = (); @@ -181,7 +183,7 @@ my $ISACMD=""; if($HIP_PLATFORM eq "hcc"){ $ISACMD .= "$HIP_PATH/bin/hipgenisa.sh "; $ISACMD .= $ROCM_PATH; - if($ARGV[0] eq "--gencodeobject"){ + if($ARGV[0] eq "--genco"){ foreach $isaarg (@ARGV[1..$#ARGV]){ $ISACMD .= " "; $ISACMD .= $isaarg; @@ -196,7 +198,7 @@ if($HIP_PLATFORM eq "hcc"){ if($HIP_PLATFORM eq "nvcc"){ $ISACMD .= "$HIP_PATH/bin/hipcc -ptx "; - if($ARGV[0] eq "--gencodeobject"){ + if($ARGV[0] eq "--genco"){ foreach $isaarg (@ARGV[1..$#ARGV]){ $ISACMD .= " "; $ISACMD .= $isaarg; @@ -240,6 +242,10 @@ foreach $arg (@ARGV) $printHipVersion = 1; $runCmd = 0; } + if($arg eq '-M') { + $compileOnly = 1; + $buildDeps = 1; + } if ($arg =~ m/^-/) { # options start with - @@ -278,6 +284,9 @@ if ($hasC and $HIP_PLATFORM eq 'nvcc') { if ($hasCU and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -x c++"; } +if ($buildDeps and $HIP_PLATFORM eq 'nvcc') { + $HIPCXXFLAGS .= " -M -D__CUDACC__"; +} if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc') { diff --git a/projects/hip/bin/hipgenisa.sh b/projects/hip/bin/hipgenisa.sh index e60abbf78f..f52b1c5f4e 100755 --- a/projects/hip/bin/hipgenisa.sh +++ b/projects/hip/bin/hipgenisa.sh @@ -21,11 +21,22 @@ HIP_PATH="$( command cd -P "$( dirname "$SOURCE" )/.." && pwd )" export KMDUMPISA=1 export KMDUMPLLVM=1 -mkdir /tmp/hipgenisa -$HIP_PATH/bin/hipcc $FILE_NAMES -o /tmp/hipgenisa/a.out -mv dump.* /tmp/hipgenisa/ -$ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$TARGET -filetype=obj /tmp/hipgenisa/dump.isa -o /tmp/hipgenisa/dump.o -$ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa /tmp/hipgenisa/dump.o -o $OUTPUT_FILE -rm -r /tmp/hipgenisa -export KMDUMPISA=0 -export KMDUMPLLVM=0 +hipgenisa_dir=`mktemp -d --tmpdir=/tmp hip.XXXXXXXX`; +sed 's/extern \+"C" \+//g' $FILE_NAMES > $FILE_NAMES.kernel.tmp.cpp +echo " +int main(){} +" >> $FILE_NAMES.kernel.tmp.cpp +$HIP_PATH/bin/hipcc $FILE_NAMES.kernel.tmp.cpp -o $hipgenisa_dir/a.out +mv dump.* $hipgenisa_dir +$ROCM_PATH/hcc-lc/bin/llvm-mc -arch=amdgcn -mcpu=$TARGET -filetype=obj $hipgenisa_dir/dump.isa -o $hipgenisa_dir/dump.o +$ROCM_PATH/llvm/bin/clang -target amdgcn--amdhsa $hipgenisa_dir/dump.o -o $hipgenisa_dir/dump.co +map_sym="" +kernels=$(objdump -t $hipgenisa_dir/dump.co | grep grid_launch_parm | sed 's/ \+/ /g; s/\t/ /g' | cut -d" " -f6) +for mangled_sym in $kernels +do + real_sym=$(c++filt $(c++filt _$mangled_sym | cut -d: -f3 | sed 's/_functor//g') | cut -d\( -f1) + map_sym="--redefine-sym $mangled_sym=$real_sym $map_sym" +done +objcopy -F elf64-little $map_sym $hipgenisa_dir/dump.co $OUTPUT_FILE +rm $FILE_NAMES.kernel.tmp.cpp +rm -r $hipgenisa_dir diff --git a/projects/hip/clang-hipify/src/Cuda2Hip.cpp b/projects/hip/clang-hipify/src/Cuda2Hip.cpp index fde5b7e55b..8959c1bc95 100644 --- a/projects/hip/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/hip/clang-hipify/src/Cuda2Hip.cpp @@ -117,30 +117,200 @@ struct cuda2hipMap { cuda2hipRename["cublas_v2.h"] = {"hipblas.h", CONV_INCLUDE, API_BLAS}; // Error codes and return types + cuda2hipRename["CUresult"] = {"hipError_t", CONV_TYPE, API_DRIVER}; cuda2hipRename["cudaError_t"] = {"hipError_t", CONV_TYPE, API_RUNTIME}; cuda2hipRename["cudaError"] = {"hipError", CONV_TYPE, API_RUNTIME}; - cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorMemoryAllocation"] = {"hipErrorMemoryAllocation", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorMemoryFree"] = {"hipErrorMemoryFree", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorUnknownSymbol"] = {"hipErrorUnknownSymbol", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorOutOfResources"] = {"hipErrorOutOfResources", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorInvalidValue"] = {"hipErrorInvalidValue", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorInvalidResourceHandle"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorInvalidDevice"] = {"hipErrorInvalidDevice", CONV_ERR, API_RUNTIME}; + + // CUDA Driver API error code only + cuda2hipRename["CUDA_ERROR_INVALID_CONTEXT"] = {"hipErrorInvalidContext", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_CONTEXT_ALREADY_CURRENT"] = {"hipErrorContextAlreadyCurrent", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_MAP_FAILED"] = {"hipErrorMapFailed", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_UNMAP_FAILED"] = {"hipErrorUnmapFailed", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_ARRAY_IS_MAPPED"] = {"hipErrorArrayIsMapped", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_ALREADY_MAPPED"] = {"hipErrorAlreadyMapped", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_ALREADY_ACQUIRED"] = {"hipErrorAlreadyAcquired", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_NOT_MAPPED"] = {"hipErrorNotMapped", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_NOT_MAPPED_AS_ARRAY"] = {"hipErrorNotMappedAsArray", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_NOT_MAPPED_AS_POINTER"] = {"hipErrorNotMappedAsPointer", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_CONTEXT_ALREADY_IN_USE"] = {"hipErrorContextAlreadyInUse", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_INVALID_SOURCE"] = {"hipErrorInvalidSource", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_FILE_NOT_FOUND"] = {"hipErrorFileNotFound", CONV_ERR, API_DRIVER}; + cuda2hipRename["CUDA_ERROR_NOT_FOUND"] = {"hipErrorNotFound", CONV_ERR, API_DRIVER}; + + // CUDA RT API error code only + cuda2hipRename["cudaErrorInvalidDeviceFunction"] = {"hipErrorInvalidDeviceFunction", CONV_ERR, API_RUNTIME}; + cuda2hipRename["cudaErrorInvalidConfiguration"] = {"hipErrorInvalidConfiguration", CONV_ERR, API_RUNTIME}; + cuda2hipRename["cudaErrorPriorLaunchFailure"] = {"hipErrorPriorLaunchFailure", CONV_ERR, API_RUNTIME}; cuda2hipRename["cudaErrorInvalidMemcpyDirection"] = {"hipErrorInvalidMemcpyDirection", CONV_ERR, API_RUNTIME}; cuda2hipRename["cudaErrorInvalidDevicePointer"] = {"hipErrorInvalidDevicePointer", CONV_ERR, API_RUNTIME}; - cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorInvalidDevicePointer", CONV_ERR, API_RUNTIME}; + cuda2hipRename["cudaErrorMissingConfiguration"] = {"hipErrorMissingConfiguration", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_SUCCESS"] = {"hipSuccess", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_UNKNOWN"] = {"hipErrorUnknown", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorUnknown"] = {"hipErrorUnknown", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_NOT_INITIALIZED"] = {"hipErrorNotInitialized", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorNotInitialized", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_DEINITIALIZED"] = {"hipErrorDeinitialized", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorCudartUnloading"] = {"hipErrorDeinitialized", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_OUT_OF_MEMORY"] = {"hipErrorMemoryAllocation", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorMemoryAllocation"] = {"hipErrorMemoryAllocation", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_HANDLE"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidResourceHandle"] = {"hipErrorInvalidResourceHandle", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_VALUE"] = {"hipErrorInvalidValue", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidValue"] = {"hipErrorInvalidValue", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_DEVICE"] = {"hipErrorInvalidDevice", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidDevice"] = {"hipErrorInvalidDevice", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_NOT_INITIALIZED"] = {"hipErrorInitializationError", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInitializationError"] = {"hipErrorInitializationError", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_NO_DEVICE"] = {"hipErrorNoDevice", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorNoDevice"] = {"hipErrorNoDevice", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_NOT_READY"] = {"hipErrorNotReady", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorNotReady"] = {"hipErrorNotReady", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PEER_ACCESS_NOT_ENABLED"] = {"hipErrorPeerAccessNotEnabled", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorPeerAccessNotEnabled"] = {"hipErrorPeerAccessNotEnabled", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED"] = {"hipErrorPeerAccessAlreadyEnabled", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorPeerAccessAlreadyEnabled"] = {"hipErrorPeerAccessAlreadyEnabled", CONV_ERR, API_RUNTIME}; - // NOTE: no corresponding error type in CUDA - //cuda2hipRename["cudaErrorRuntimeMemory"] = {"hipErrorRuntimeMemory", CONV_ERR, API_RUNTIME}; - //cuda2hipRename["cudaErrorRuntimeOther"] = {"hipErrorRuntimeOther", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PEER_ACCESS_UNSUPPORTED"] = {"hipErrorPeerAccessUnsupported", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorPeerAccessUnsupported"] = {"hipErrorPeerAccessUnsupported", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_PTX"] = {"hipErrorInvalidKernelFile", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidPtx"] = {"hipErrorInvalidKernelFile", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_GRAPHICS_CONTEXT"] = {"hipErrorInvalidGraphicsContext", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidGraphicsContext"] = {"hipErrorInvalidGraphicsContext", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND"] = {"hipErrorSharedObjectSymbolNotFound", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorSharedObjectSymbolNotFound"] = {"hipErrorSharedObjectSymbolNotFound", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_SHARED_OBJECT_INIT_FAILED"] = {"hipErrorSharedObjectInitFailed", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorSharedObjectInitFailed"] = {"hipErrorSharedObjectInitFailed", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_OPERATING_SYSTEM"] = {"hipErrorOperatingSystem", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorOperatingSystem"] = {"hipErrorOperatingSystem", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_ILLEGAL_ADDRESS"] = {"hipErrorIllegalAddress", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorIllegalAddress"] = {"hipErrorIllegalAddress", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_LAUNCH_FAILED"] = {"hipErrorLaunchFailure", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorLaunchFailure"] = {"hipErrorLaunchFailure", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_LAUNCH_TIMEOUT"] = {"hipErrorLaunchTimeOut", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorLaunchTimeout"] = {"hipErrorLaunchTimeOut", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES"] = {"hipErrorLaunchOutOfResources", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorLaunchOutOfResources"] = {"hipErrorLaunchOutOfResources", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_ECC_UNCORRECTABLE"] = {"hipErrorECCNotCorrectable", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorECCUncorrectable"] = {"hipErrorECCNotCorrectable", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED"] = {"hipErrorHostMemoryAlreadyRegistered", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorHostMemoryAlreadyRegistered"] = {"hipErrorHostMemoryAlreadyRegistered", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED"] = {"hipErrorHostMemoryNotRegistered", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorHostMemoryNotRegistered"] = {"hipErrorHostMemoryNotRegistered", CONV_ERR, API_RUNTIME}; + cuda2hipRename["CUDA_ERROR_NO_BINARY_FOR_GPU"] = {"hipErrorNoBinaryForGpu", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorNoKernelImageForDevice"] = {"hipErrorNoBinaryForGpu", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_UNSUPPORTED_LIMIT"] = {"hipErrorUnsupportedLimit", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorUnsupportedLimit"] = {"hipErrorUnsupportedLimit", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_INVALID_IMAGE"] = {"hipErrorInvalidImage", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorInvalidKernelImage"] = {"hipErrorInvalidImage", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PROFILER_DISABLED"] = {"hipErrorProfilerDisabled", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorProfilerDisabled"] = {"hipErrorProfilerDisabled", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PROFILER_NOT_INITIALIZED"] = {"hipErrorProfilerNotInitialized", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorProfilerNotInitialized"] = {"hipErrorProfilerNotInitialized", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PROFILER_ALREADY_STARTED"] = {"hipErrorProfilerAlreadyStarted", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorProfilerAlreadyStarted"] = {"hipErrorProfilerAlreadyStarted", CONV_ERR, API_RUNTIME}; + + cuda2hipRename["CUDA_ERROR_PROFILER_ALREADY_STOPPED"] = {"hipErrorProfilerAlreadyStopped", CONV_ERR, API_DRIVER}; + cuda2hipRename["cudaErrorProfilerAlreadyStopped"] = {"hipErrorProfilerAlreadyStopped", CONV_ERR, API_RUNTIME}; + + ///////////////////////////// CUDA DRIVER API ///////////////////////////// + // Types + // NOTE: CUdevice might be changed to typedef int in the future. + cuda2hipRename["CUdevice"] = {"hipDevice_t", CONV_TYPE, API_DRIVER}; + + cuda2hipRename["CUdevice_attribute_enum"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUdevice_attribute"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER}; + + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK"] = {"hipDeviceAttributeMaxThreadsPerBlock", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X"] = {"hipDeviceAttributeMaxBlockDimX", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y"] = {"hipDeviceAttributeMaxBlockDimY", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z"] = {"hipDeviceAttributeMaxBlockDimZ", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X"] = {"hipDeviceAttributeMaxGridDimX", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y"] = {"hipDeviceAttributeMaxGridDimY", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z"] = {"hipDeviceAttributeMaxGridDimZ", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK"] = {"hipDeviceAttributeMaxSharedMemoryPerBlock", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY"] = {"hipDeviceAttributeTotalConstantMemory", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_WARP_SIZE"] = {"hipDeviceAttributeWarpSize", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK"] = {"hipDeviceAttributeMaxRegistersPerBlock", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_CLOCK_RATE"] = {"hipDeviceAttributeClockRate", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE"] = {"hipDeviceAttributeMemoryClockRate", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH"] = {"hipDeviceAttributeMemoryBusWidth", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH"] = {"hipDeviceAttributeMultiprocessorCount", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_MODE"] = {"hipDeviceAttributeComputeMode", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE"] = {"hipDeviceAttributeL2CacheSize", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR"] = {"hipDeviceAttributeMaxThreadsPerMultiProcessor", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR"] = {"hipDeviceAttributeComputeCapabilityMajor", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR"] = {"hipDeviceAttributeComputeCapabilityMinor", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS"] = {"hipDeviceAttributeConcurrentKernels", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_PCI_BUS_ID"] = {"hipDeviceAttributePciBusId", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID"] = {"hipDeviceAttributePciDeviceId", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR"] = {"hipDeviceAttributeMaxSharedMemoryPerMultiprocessor", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD"] = {"hipDeviceAttributeIsMultiGpuBoard", CONV_DEV, API_DRIVER}; + + cuda2hipRename["CUdevprop_st"] = {"hipDeviceProp_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUdevprop"] = {"hipDeviceProp_t", CONV_TYPE, API_DRIVER}; + + // TODO: Analogues enum is needed in HIP. Couldn't map enum to struct hipPointerAttribute_t. + // TODO: Do for Pointer Attributes the same as for Device Attributes. + // cuda2hipRename["CUpointer_attribute_enum"] = {"hipPointerAttribute_t", CONV_TYPE, API_DRIVER}; + // cuda2hipRename["CUpointer_attribute"] = {"hipPointerAttribute_t", CONV_TYPE, API_DRIVER}; + + cuda2hipRename["CUfunction"] = {"hipFunction_t", CONV_TYPE, API_DRIVER}; + + // unsupported yet by HIP + // cuda2hipRename["CUfunction_attribute_enum"] = {"hipFuncAttribute_t", CONV_TYPE, API_DRIVER}; + // cuda2hipRename["CUfunction_attribute"] = {"hipFuncAttribute_t", CONV_TYPE, API_DRIVER}; + + cuda2hipRename["CUfunc_cache_enum"] = {"hipFuncCache", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUfunc_cache"] = {"hipFuncCache", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CU_FUNC_CACHE_PREFER_NONE"] = {"hipFuncCachePreferNone", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_FUNC_CACHE_PREFER_SHARED"] = {"hipFuncCachePreferShared", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_FUNC_CACHE_PREFER_L1"] = {"hipFuncCachePreferL1", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_FUNC_CACHE_PREFER_EQUAL"] = {"hipFuncCachePreferEqual", CONV_DEV, API_DRIVER}; + + cuda2hipRename["CUsharedconfig_enum"] = {"hipSharedMemConfig", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUsharedconfig"] = {"hipSharedMemConfig", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE"] = {"hipSharedMemBankSizeDefault", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE"] = {"hipSharedMemBankSizeFourByte", CONV_DEV, API_DRIVER}; + cuda2hipRename["CU_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE"] = {"hipSharedMemBankSizeEightByte", CONV_DEV, API_DRIVER}; + + cuda2hipRename["CUcontext"] = {"hipCtx_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUmodule"] = {"hipModule_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUevent"] = {"hipEvent_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUstream"] = {"hipStream_t", CONV_TYPE, API_DRIVER}; + + /////////////////////////////// CUDA RT API /////////////////////////////// // Error API cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR, API_RUNTIME}; cuda2hipRename["cudaPeekAtLastError"] = {"hipPeekAtLastError", CONV_ERR, API_RUNTIME}; @@ -279,6 +449,7 @@ struct cuda2hipMap { cuda2hipRename["cudaDevAttrL2CacheSize"] = {"hipDeviceAttributeL2CacheSize", CONV_DEV, API_RUNTIME}; cuda2hipRename["cudaDevAttrMaxThreadsPerMultiProcessor"] = {"hipDeviceAttributeMaxThreadsPerMultiProcessor", CONV_DEV, API_RUNTIME}; cuda2hipRename["cudaDevAttrComputeCapabilityMajor"] = {"hipDeviceAttributeComputeCapabilityMajor", CONV_DEV, API_RUNTIME}; + cuda2hipRename["cudaDevAttrComputeCapabilityMinor"] = {"hipDeviceAttributeComputeCapabilityMinor", CONV_DEV, API_RUNTIME}; cuda2hipRename["cudaDevAttrConcurrentKernels"] = {"hipDeviceAttributeConcurrentKernels", CONV_DEV, API_RUNTIME}; cuda2hipRename["cudaDevAttrPciBusId"] = {"hipDeviceAttributePciBusId", CONV_DEV, API_RUNTIME}; cuda2hipRename["cudaDevAttrPciDeviceId"] = {"hipDeviceAttributePciDeviceId", CONV_DEV, API_RUNTIME}; @@ -1038,8 +1209,7 @@ static void processString(StringRef s, const cuda2hipMap &map, int64_t countReps[CONV_LAST], int64_t countApiReps[API_LAST]) { size_t begin = 0; - while ((begin = s.find("cuda", begin)) != StringRef::npos || - (begin = s.find("cublas", begin)) != StringRef::npos) { + while ((begin = s.find("cu", begin)) != StringRef::npos) { const size_t end = s.find_first_of(" ", begin + 4); StringRef name = s.slice(begin, end); const auto found = map.cuda2hipRename.find(name); @@ -1709,7 +1879,7 @@ static cl::opt void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callback) { Finder.addMatcher(callExpr(isExpansionInMainFile(), - callee(functionDecl(matchesName("cuda.*|cublas.*")))) + callee(functionDecl(matchesName("cu.*")))) .bind("cudaCall"), Callback); Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), Callback); @@ -1720,7 +1890,7 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac Callback); Finder.addMatcher(declRefExpr(isExpansionInMainFile(), to(enumConstantDecl( - matchesName("cuda.*|cublas.*|CUDA.*|CUBLAS*")))) + matchesName("cu.*|CU.*")))) .bind("cudaEnumConstantRef"), Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), @@ -1728,36 +1898,36 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac .bind("cudaEnumConstantDecl"), Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), - hasType(typedefDecl(matchesName("cuda.*|cublas.*")))) + hasType(typedefDecl(matchesName("cu.*|CU.*")))) .bind("cudaTypedefVar"), Callback); // Array of elements of typedef type, Example: cudaStream_t streams[2]; Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(arrayType(hasElementType(typedefType( - hasDeclaration(typedefDecl(matchesName("cuda.*|cublas.*")))))))) + hasDeclaration(typedefDecl(matchesName("cu.*|CU.*")))))))) .bind("cudaTypedefVar"), Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), - hasType(cxxRecordDecl(matchesName("cuda.*|cublas.*")))) + hasType(cxxRecordDecl(matchesName("cu.*|CU.*")))) .bind("cudaStructVar"), Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(pointsTo(cxxRecordDecl( - matchesName("cuda.*|cublas.*"))))) + matchesName("cu.*|CU.*"))))) .bind("cudaStructVarPtr"), Callback); Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), - hasType(namedDecl(matchesName("cuda.*|cublas.*")))) + hasType(namedDecl(matchesName("cu.*|CU.*")))) .bind("cudaParamDecl"), Callback); Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(pointsTo(namedDecl( - matchesName("cuda.*|cublas.*"))))) + matchesName("cu.*|CU.*"))))) .bind("cudaParamDeclPtr"), Callback); Finder.addMatcher(expr(isExpansionInMainFile(), sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration( - cxxRecordDecl(matchesName("cuda.*|cublas.*"))))))) + cxxRecordDecl(matchesName("cu.*|CU.*"))))))) .bind("cudaStructSizeOf"), Callback); Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), diff --git a/projects/hip/cmake/CMakeLists.txt b/projects/hip/cmake/CMakeLists.txt deleted file mode 100644 index ff5728db4f..0000000000 --- a/projects/hip/cmake/CMakeLists.txt +++ /dev/null @@ -1,12 +0,0 @@ - - -#Our FindHCC is in the same directory as the this CMakeLists.txt, which is CMAKE_SOURCE_DIR -# Let cmake know where to find FindHCC.cmake -# - -set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_SOURCE_DIR}) - -find_package(HIP REQUIRED) - -message(STATUS "HIP_PLATFORM = ${HIP_PLATFORM}") - diff --git a/projects/hip/cmake/FindHCC.cmake b/projects/hip/cmake/FindHCC.cmake deleted file mode 100644 index 98f9d94f31..0000000000 --- a/projects/hip/cmake/FindHCC.cmake +++ /dev/null @@ -1,73 +0,0 @@ - -# findHCC does not currently address versioning, i.e. -# a rich directory structure where version number is a subdirectory under root -# Also, supported only on UNIX 64 bit systems. - -if(UNIX) - if(CMAKE_SIZEOF_VOID_P EQUAL 8) - - find_library(HSA_LIBRARY - NAMES hsa-runtime64 - PATHS - ENV HSA_PATH - /opt/rocm/hsa - PATH_SUFFIXES - lib) - - if( NOT DEFINED ENV{HSA_PATH} ) - set( ENV{HSA_PATH} /opt/rocm/hsa) - endif() - - find_program(HCC - NAMES hcc - PATHS - ENV HCC_PATH - /opt/rocm/hcc - PATH_SUFFIXES - /bin) - - if( NOT DEFINED ENV{HCC_PATH} ) - set( ENV{HCC_PATH} /opt/rocm/hcc) - endif() - -# this is now dynamic -# find_library(AMP_LIBRARY -# NAMES mcwamp -# PATHS -# ENV NCC_PATH -# /opt/rocm/hcc -# PATH_SUFFIXES -# /lib) - - find_path(HCC_INCLUDE_DIR - NAMES - hc.hpp - PATHS - ENV NCC_PATH - /opt/rocm/hcc - PATH_SUFFIXES - /include) - - - - set(HSA_LIBRARIES ${HSA_LIBRARY}) - #set(HCC_LIBRARIES ${AMP_LIBRARY}) - set(HCC_INCLUDE_DIRS ${HCC_INCLUDE_DIR}) - - include(FindPackageHandleStandardArgs) - find_package_handle_standard_args( - HCC - FOUND_VAR HCC_FOUND - REQUIRED_VARS HSA_LIBRARIES HCC_INCLUDE_DIRS HCC) - - mark_as_advanced( - HSA_LIBRARIES - HCC_INCLUDE_DIRS - ) - - else() - message(SEND_ERROR "HCC is currently supported only on 64 bit UNIX platforms") - endif() -else() - message(SEND_ERROR "HCC is currently supported on unix platforms") -endif() diff --git a/projects/hip/cmake/FindHIP.cmake b/projects/hip/cmake/FindHIP.cmake index c0e3ac5d6f..7b129d1550 100644 --- a/projects/hip/cmake/FindHIP.cmake +++ b/projects/hip/cmake/FindHIP.cmake @@ -1,56 +1,463 @@ +############################################################################### +# FindHIP.cmake +############################################################################### -find_package(CUDA) +############################################################################### +# SET: Variable defaults +############################################################################### +# User defined flags +set(HIP_HIPCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HIPCC") +set(HIP_HCC_FLAGS "" CACHE STRING "Semicolon delimited flags for HCC") +set(HIP_NVCC_FLAGS "" CACHE STRING "Semicolon delimted flags for NVCC") +mark_as_advanced(HIP_HIPCC_FLAGS HIP_HCC_FLAGS HIP_NVCC_FLAGS) +set(_hip_configuration_types ${CMAKE_CONFIGURATION_TYPES} ${CMAKE_BUILD_TYPE} Debug MinSizeRel Release RelWithDebInfo) +list(REMOVE_DUPLICATES _hip_configuration_types) +foreach(config ${_hip_configuration_types}) + string(TOUPPER ${config} config_upper) + set(HIP_HIPCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HIPCC") + set(HIP_HCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for HCC") + set(HIP_NVCC_FLAGS_${config_upper} "" CACHE STRING "Semicolon delimited flags for NVCC") + mark_as_advanced(HIP_HIPCC_FLAGS_${config_upper} HIP_HCC_FLAGS_${config_upper} HIP_NVCC_FLAGS_${config_upper}) +endforeach() +option(HIP_HOST_COMPILATION_CPP "Host code compilation mode" ON) +mark_as_advanced(HIP_HOST_COMPILATION_CPP) -if( CUDA_INCLUDE_DIRS AND CUDA_VERSION AND CUDA_NVCC_EXECUTABLE) - message(STATUS "CUDA_VERSION = ${CUDA_VERSION}") - message(STATUS "CUDA_INCLUDE_DIRS = ${CUDA_INCLUDE_DIRS}") - message(STATUS "CUDA_NVCC_EXECUTABLE = ${CUDA_NVCC_EXECUTABLE}") - - set( HIP_PLATFORM "nvcc" ) - - #export the environment variable, so that HIPCC can find it. - set(ENV{HIP_PLATFORM} nvcc) - -else() - find_package(HCC) - - message(STATUS "HCC_FOUND = ${HCC_FOUND}") - message(STATUS "HCC = ${HCC}") - message(STATUS "HCC_INCLUDE_DIRS = ${HCC_INCLUDE_DIRS}") - message(STATUS "HSA_LIBRARIES = ${HSA_LIBRARIES}") +############################################################################### +# FIND: HIP and associated helper binaries +############################################################################### +# HIP is supported on Linux only +if(UNIX AND NOT APPLE AND NOT CYGWIN) + # Search for HIP installation + if(NOT HIP_ROOT_DIR) + # Search in user specified path first + find_path( + HIP_ROOT_DIR + NAMES hipconfig + PATHS + ENV ROCM_PATH + ENV HIP_PATH + PATH_SUFFIXES bin + DOC "HIP installed location" + NO_DEFAULT_PATH + ) + # Now search in default path + find_path( + HIP_ROOT_DIR + NAMES hipconfig + PATHS + /opt/rocm + /opt/rocm/hip + PATH_SUFFIXES bin + DOC "HIP installed location" + ) - if( ${HCC_FOUND} STREQUAL "TRUE" ) - - # This directory is hip/cmake! HIP_PATH should be one directory up! - set (HIP_PATH $ENV{HIP_PATH}) - if (NOT DEFINED HIP_PATH) - set (HIP_PATH ${CMAKE_CURRENT_SOURCE_DIR}/..) - set( ENV{HIP_PATH} ${HIP_PATH}) - endif() + # Check if we found HIP installation + if(HIP_ROOT_DIR) + # If so, fix the path + string(REGEX REPLACE "[/\\\\]?bin[64]*[/\\\\]?$" "" HIP_ROOT_DIR ${HIP_ROOT_DIR}) + # And push it back to the cache + set(HIP_ROOT_DIR ${HIP_ROOT_DIR} CACHE PATH "HIP installed location" FORCE) + endif() + if(NOT EXISTS ${HIP_ROOT_DIR}) + if(HIP_FIND_REQUIRED) + message(FATAL_ERROR "Specify HIP_ROOT_DIR") + elseif(NOT HIP_FIND_QUIETLY) + message("HIP_ROOT_DIR not found or specified") + endif() + endif() + endif() - message(STATUS "ENV HIP_PATH = $ENV{HIP_PATH}") - - find_program(HIPCC - NAMES hipcc - PATHS - ENV HIP_PATH - PATH_SUFFIXES - /bin) - - message(STATUS "HIPCC = ${HIPCC}") - - if( DEFINED HIPCC) - - set( HIP_PLATFORM "hcc" ) - #export the environment variable, so that HIPCC can find it. - set(ENV{HIP_PLATFORM} "hcc") - set (CMAKE_CXX_COMPILER ${HIPCC}) - - else() - message(SEND_ERROR "Did not find HIPCC") - endif() - else() - message(SEND_ERROR "hcc not found") - endif() - + # Find HIPCC executable + find_program( + HIP_HIPCC_EXECUTABLE + NAMES hipcc + PATHS + "${HIP_ROOT_DIR}" + ENV ROCM_PATH + ENV HIP_PATH + /opt/rocm + /opt/rocm/hip + PATH_SUFFIXES bin + NO_DEFAULT_PATH + ) + if(NOT HIP_HIPCC_EXECUTABLE) + # Now search in default paths + find_program(HIP_HIPCC_EXECUTABLE hipcc) + endif() + mark_as_advanced(HIP_HIPCC_EXECUTABLE) + + # Find HIPCONFIG executable + find_program( + HIP_HIPCONFIG_EXECUTABLE + NAMES hipconfig + PATHS + "${HIP_ROOT_DIR}" + ENV ROCM_PATH + ENV HIP_PATH + /opt/rocm + /opt/rocm/hip + PATH_SUFFIXES bin + NO_DEFAULT_PATH + ) + if(NOT HIP_HIPCONFIG_EXECUTABLE) + # Now search in default paths + find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig) + endif() + mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE) + + if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_VERSION) + # Compute the version + execute_process( + COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --version + OUTPUT_VARIABLE _hip_version + ERROR_VARIABLE _hip_error + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_STRIP_TRAILING_WHITESPACE + ) + if(NOT _hip_error) + set(HIP_VERSION ${_hip_version} CACHE STRING "Version of HIP as computed from hipcc") + else() + set(HIP_VERSION "0.0.0" CACHE STRING "Version of HIP as computed by FindHIP()") + endif() + mark_as_advanced(HIP_VERSION) + endif() + if(HIP_VERSION) + string(REPLACE "." ";" _hip_version_list "${HIP_VERSION}") + list(GET _hip_version_list 0 HIP_VERSION_MAJOR) + list(GET _hip_version_list 1 HIP_VERSION_MINOR) + list(GET _hip_version_list 2 HIP_VERSION_PATCH) + set(HIP_VERSION_STRING "${HIP_VERSION}") + endif() + + if(HIP_HIPCONFIG_EXECUTABLE AND NOT HIP_PLATFORM) + # Compute the platform + execute_process( + COMMAND ${HIP_HIPCONFIG_EXECUTABLE} --platform + OUTPUT_VARIABLE _hip_platform + OUTPUT_STRIP_TRAILING_WHITESPACE + ) + set(HIP_PLATFORM ${_hip_platform} CACHE STRING "HIP platform as computed by hipconfig") + mark_as_advanced(HIP_PLATFORM) + endif() endif() + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args( + HIP + REQUIRED_VARS + HIP_ROOT_DIR + HIP_HIPCC_EXECUTABLE + HIP_HIPCONFIG_EXECUTABLE + HIP_PLATFORM + VERSION_VAR HIP_VERSION + ) + +############################################################################### +# MACRO: Locate helper files +############################################################################### +macro(HIP_FIND_HELPER_FILE _name _extension) + set(_hip_full_name "${_name}.${_extension}") + get_filename_component(CMAKE_CURRENT_LIST_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH) + set(HIP_${_name} "${CMAKE_CURRENT_LIST_DIR}/FindHIP/${_hip_full_name}") + if(NOT EXISTS "${HIP_${_name}}") + set(error_message "${_hip_full_name} not found in ${CMAKE_CURRENT_LIST_DIR}/FindHIP") + if(HIP_FIND_REQUIRED) + message(FATAL_ERROR "${error_message}") + else() + if(NOT HIP_FIND_QUIETLY) + message(STATUS "${error_message}") + endif() + endif() + endif() + # Set this variable as internal, so the user isn't bugged with it. + set(HIP_${_name} ${HIP_${_name}} CACHE INTERNAL "Location of ${_full_name}" FORCE) +endmacro() + +############################################################################### +hip_find_helper_file(run_make2cmake cmake) +hip_find_helper_file(run_hipcc cmake) +############################################################################### + +############################################################################### +# MACRO: Seperate the options from the sources +############################################################################### +macro(HIP_GET_SOURCES_AND_OPTIONS _sources _cmake_options _hipcc_options _hcc_options _nvcc_options) + set(${_sources}) + set(${_cmake_options}) + set(${_hipcc_options}) + set(${_hcc_options}) + set(${_nvcc_options}) + set(_hipcc_found_options FALSE) + set(_hcc_found_options FALSE) + set(_nvcc_found_options FALSE) + foreach(arg ${ARGN}) + if("x${arg}" STREQUAL "xHIPCC_OPTIONS") + set(_hipcc_found_options TRUE) + elseif("x${arg}" STREQUAL "xHCC_OPTIONS") + set(_hcc_found_options TRUE) + elseif("x${arg}" STREQUAL "xNVCC_OPTIONS") + set(_nvcc_found_options TRUE) + elseif( + "x${arg}" STREQUAL "xEXCLUDE_FROM_ALL" OR + "x${arg}" STREQUAL "xSTATIC" OR + "x${arg}" STREQUAL "xSHARED" OR + "x${arg}" STREQUAL "xMODULE" + ) + list(APPEND ${_cmake_options} ${arg}) + else() + if(_hipcc_found_options) + list(APPEND ${_hipcc_options} ${arg}) + elseif(_hcc_found_options) + list(APPEND ${_hcc_options} ${arg}) + elseif(_nvcc_found_options) + list(APPEND ${_nvcc_options} ${arg}) + else() + # Assume this is a file + list(APPEND ${_sources} ${arg}) + endif() + endif() + endforeach() +endmacro() + +############################################################################### +# MACRO: Add include directories to pass to the hipcc command +############################################################################### +set(HIP_HIPCC_INCLUDE_ARGS_USER "") +macro(HIP_INCLUDE_DIRECTORIES) + foreach(dir ${ARGN}) + list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER -I${dir}) + endforeach() +endmacro() + +############################################################################### +# FUNCTION: Helper to avoid clashes of files with the same basename but different paths +############################################################################### +function(HIP_COMPUTE_BUILD_PATH path build_path) + # Convert to cmake style paths + file(TO_CMAKE_PATH "${path}" bpath) + if (IS_ABSOLUTE "${bpath}") + string(FIND "${bpath}" "${CMAKE_CURRENT_BINARY_DIR}" _binary_dir_pos) + if (_binary_dir_pos EQUAL 0) + file(RELATIVE_PATH bpath "${CMAKE_CURRENT_BINARY_DIR}" "${bpath}") + else() + file(RELATIVE_PATH bpath "${CMAKE_CURRENT_SOURCE_DIR}" "${bpath}") + endif() + endif() + + # Remove leading / + string(REGEX REPLACE "^[/]+" "" bpath "${bpath}") + # Avoid absolute paths by removing ':' + string(REPLACE ":" "_" bpath "${bpath}") + # Avoid relative paths that go up the tree + string(REPLACE "../" "__/" bpath "${bpath}") + # Avoid spaces + string(REPLACE " " "_" bpath "${bpath}") + # Strip off the filename + get_filename_component(bpath "${bpath}" PATH) + + set(${build_path} "${bpath}" PARENT_SCOPE) +endfunction() + +############################################################################### +# MACRO: Parse OPTIONS from ARGN & set variables prefixed by _option_prefix +############################################################################### +macro(HIP_PARSE_HIPCC_OPTIONS _option_prefix) + set(_hip_found_config) + foreach(arg ${ARGN}) + # Determine if we are dealing with a per-configuration flag + foreach(config ${_hip_configuration_types}) + string(TOUPPER ${config} config_upper) + if(arg STREQUAL "${config_upper}") + set(_hip_found_config _${arg}) + # Clear arg to prevent it from being processed anymore + set(arg) + endif() + endforeach() + if(arg) + list(APPEND ${_option_prefix}${_hip_found_config} "${arg}") + endif() + endforeach() +endmacro() + +############################################################################### +# MACRO: Try and include dependency file if it exists +############################################################################### +macro(HIP_INCLUDE_HIPCC_DEPENDENCIES dependency_file) + set(HIP_HIPCC_DEPEND) + set(HIP_HIPCC_DEPEND_REGENERATE FALSE) + + # Create the dependency file if it doesn't exist + if(NOT EXISTS ${dependency_file}) + file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n") + endif() + # Include the dependency file + include(${dependency_file}) + + # Verify the existence of all the included files + if(HIP_HIPCC_DEPEND) + foreach(f ${HIP_HIPCC_DEPEND}) + if(NOT EXISTS ${f}) + # If they aren't there, regenerate the file again + set(HIP_HIPCC_DEPEND_REGENERATE TRUE) + endif() + endforeach() + else() + # No dependencies, so regenerate the file + set(CUDA_NVCC_DEPEND_REGENERATE TRUE) + endif() + + # Regenerate the dependency file if needed + if(HIP_HIPCC_DEPEND_REGENERATE) + set(HIP_HIPCC_DEPEND ${dependency_file}) + file(WRITE ${dependency_file} "# Generated by: FindHIP.cmake. Do not edit.\n") + endif() +endmacro() + +############################################################################### +# MACRO: Prepare cmake commands for the target +############################################################################### +macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files) + set(_hip_flags "") + set(_hip_build_configuration "${CMAKE_BUILD_TYPE}") + if(HIP_HOST_COMPILATION_CPP) + set(HIP_C_OR_CXX CXX) + else() + set(HIP_C_OR_CXX C) + endif() + set(generated_extension ${CMAKE_${HIP_C_OR_CXX}_OUTPUT_EXTENSION}) + + # Initialize list of includes with those specified by the user. Append with + # ones specified to cmake directly. + set(HIP_HIPCC_INCLUDE_ARGS ${HIP_HIPCC_INCLUDE_ARGS_USER}) + get_directory_property(_hip_include_directories INCLUDE_DIRECTORIES) + list(REMOVE_DUPLICATES _hip_include_directories) + if(_hip_include_directories) + foreach(dir ${_hip_include_directories}) + list(APPEND HIP_HIPCC_INCLUDE_ARGS -I${dir}) + endforeach() + endif() + + HIP_GET_SOURCES_AND_OPTIONS(_hip_sources _hip_cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) + HIP_PARSE_HIPCC_OPTIONS(HIP_HIPCC_FLAGS ${_hipcc_options}) + HIP_PARSE_HIPCC_OPTIONS(HIP_HCC_FLAGS ${_hcc_options}) + HIP_PARSE_HIPCC_OPTIONS(HIP_NVCC_FLAGS ${_nvcc_options}) + + # Set host compiler + set(HIP_HOST_COMPILER "${CMAKE_${HIP_C_OR_CXX}_COMPILER}") + + # Set compiler flags + set(_HIP_HOST_FLAGS "set(CMAKE_HOST_FLAGS ${CMAKE_${HIP_C_OR_CXX}_FLAGS})") + set(_HIP_HIPCC_FLAGS "set(HIP_HIPCC_FLAGS ${HIP_HIPCC_FLAGS})") + set(_HIP_HCC_FLAGS "set(HIP_HCC_FLAGS ${HIP_HCC_FLAGS})") + set(_HIP_NVCC_FLAGS "set(HIP_NVCC_FLAGS ${HIP_NVCC_FLAGS})") + foreach(config ${_hip_configuration_types}) + string(TOUPPER ${config} config_upper) + set(_HIP_HOST_FLAGS "${_HIP_HOST_FLAGS}\nset(CMAKE_HOST_FLAGS_${config_upper} ${CMAKE_${HIP_C_OR_CXX}_FLAGS_${config_upper}})") + set(_HIP_HIPCC_FLAGS "${_HIP_HIPCC_FLAGS}\nset(HIP_HIPCC_FLAGS_${config_upper} ${HIP_HIPCC_FLAGS_${config_upper}})") + set(_HIP_HCC_FLAGS "${_HIP_HCC_FLAGS}\nset(HIP_HCC_FLAGS_${config_upper} ${HIP_HCC_FLAGS_${config_upper}})") + set(_HIP_NVCC_FLAGS "${_HIP_NVCC_FLAGS}\nset(HIP_NVCC_FLAGS_${config_upper} ${HIP_NVCC_FLAGS_${config_upper}})") + endforeach() + + # Reset the output variable + set(_hip_generated_files "") + + # Iterate over all arguments and create custom commands for all source files + foreach(file ${ARGN}) + # Ignore any file marked as a HEADER_FILE_ONLY + get_source_file_property(_is_header ${file} HEADER_FILE_ONLY) + # Allow per source file overrides of the format. Also allows compiling non .cu files. + get_source_file_property(_hip_source_format ${file} HIP_SOURCE_PROPERTY_FORMAT) + if((${file} MATCHES "\\.cu$" OR _hip_source_format) AND NOT _is_header) + set(host_flag FALSE) + else() + set(host_flag TRUE) + endif() + + if (NOT host_flag) + # Determine output directory + HIP_COMPUTE_BUILD_PATH("${file}" hip_build_path) + set(hip_compile_output_dir "${CMAKE_CURRENT_BINARY_DIR}/CMakeFiles/${_target}.dir/${hip_build_path}") + + get_filename_component(basename ${file} NAME) + set(generated_file_path "${hip_compile_output_dir}/${CMAKE_CFG_INTDIR}") + set(generated_file_basename "${_target}_generated_${basename}${generated_extension}") + + # Set file names + set(generated_file "${generated_file_path}/${generated_file_basename}") + set(cmake_dependency_file "${hip_compile_output_dir}/${generated_file_basename}.depend") + set(custom_target_script_pregen "${hip_compile_output_dir}/${generated_file_basename}.cmake.pre-gen") + set(custom_target_script "${hip_compile_output_dir}/${generated_file_basename}.cmake") + + # Set properties for object files + set_source_files_properties("${generated_file}" + PROPERTIES + EXTERNAL_OBJECT true # This is an object file not to be compiled, but only be linked + ) + + # Don't add CMAKE_CURRENT_SOURCE_DIR if the path is already an absolute path + get_filename_component(file_path "${file}" PATH) + if(IS_ABSOLUTE "${file_path}") + set(source_file "${file}") + else() + set(source_file "${CMAKE_CURRENT_SOURCE_DIR}/${file}") + endif() + + # Bring in the dependencies + HIP_INCLUDE_HIPCC_DEPENDENCIES(${cmake_dependency_file}) + + # Configure the build script + configure_file("${HIP_run_hipcc}" "${custom_target_script_pregen}" @ONLY) + file(GENERATE + OUTPUT "${custom_target_script}" + INPUT "${custom_target_script_pregen}" + ) + set(main_dep DEPENDS ${source_file}) + set(verbose_output "$(VERBOSE)") + + # Create up the comment string + file(RELATIVE_PATH generated_file_relative_path "${CMAKE_BINARY_DIR}" "${generated_file}") + set(hip_build_comment_string "Building HIPCC (${cuda_build_type}) object ${generated_file_relative_path}") + + # Build the generated file and dependency file + add_custom_command( + OUTPUT ${generated_file} + # These output files depend on the source_file and the contents of cmake_dependency_file + ${main_dep} + DEPENDS ${HIP_HIPCC_DEPEND} + DEPENDS ${custom_target_script} + # Make sure the output directory exists before trying to write to it. + COMMAND ${CMAKE_COMMAND} -E make_directory "${generated_file_path}" + COMMAND ${CMAKE_COMMAND} ARGS + -D verbose:BOOL=${verbose_output} + -D build_configuration:STRING=${_hip_build_configuration} + -D "generated_file:STRING=${generated_file}" + -P "${custom_target_script}" + WORKING_DIRECTORY "${hip_compile_output_dir}" + COMMENT "${hip_build_comment_string}" + ) + + # Make sure the build system knows the file is generated + set_source_files_properties(${generated_file} PROPERTIES GENERATED TRUE) + list(APPEND _hip_generated_files ${generated_file}) + endif() + endforeach() + + # Set the return parameter + set(${_generated_files} ${_hip_generated_files}) +endmacro() + +############################################################################### +# HIP_ADD_EXECUTABLE +############################################################################### +macro(HIP_ADD_EXECUTABLE hip_target) + # Separate the sources from the options + HIP_GET_SOURCES_AND_OPTIONS(_sources _cmake_options _hipcc_options _hcc_options _nvcc_options ${ARGN}) + HIP_PREPARE_TARGET_COMMANDS(${hip_target} OBJ _generated_files ${_sources} HIPCC_OPTIONS ${_hipcc_options} HCC_OPTIONS ${_hcc_options} NVCC_OPTIONS ${_nvcc_options}) + set(HIP_CMAKE_CXX_LINK_EXECUTABLE ${CMAKE_CXX_LINK_EXECUTABLE}) + set(CMAKE_CXX_LINK_EXECUTABLE "${HIP_HIPCC_EXECUTABLE} -o ") + add_executable(${hip_target} ${_cmake_options} ${_generated_files} ${_sources}) + set_target_properties(${hip_target} PROPERTIES LINKER_LANGUAGE ${HIP_C_OR_CXX}) + #set(CMAKE_CXX_COMPILER ${ORIGINAL_CMAKE_CXX_COMPILER}) +endmacro() + +# vim: ts=4:sw=4:expandtab:smartindent diff --git a/projects/hip/cmake/FindHIP/run_hipcc.cmake b/projects/hip/cmake/FindHIP/run_hipcc.cmake new file mode 100644 index 0000000000..8af1c72700 --- /dev/null +++ b/projects/hip/cmake/FindHIP/run_hipcc.cmake @@ -0,0 +1,163 @@ +############################################################################### +# Runs commands using HIPCC +############################################################################### + +############################################################################### +# This file runs the hipcc commands to produce the desired output file +# along with the dependency file needed by CMake to compute dependencies. +# +# Input variables: +# +# verbose:BOOL=<> OFF: Be as quiet as possible (default) +# ON : Describe each step +# build_configuration:STRING=<> Build configuration. Defaults to Debug. +# generated_file:STRING=<> File to generate. Mandatory argument. + +if(NOT build_configuration) + set(build_configuration Debug) +endif() +if(NOT generated_file) + message(FATAL_ERROR "You must specify generated_file on the command line") +endif() + +# Set these up as variables to make reading the generated file easier +set(HIP_HIPCC_EXECUTABLE "@HIP_HIPCC_EXECUTABLE@") # path +set(HIP_HOST_COMPILER "@CUDA_HOST_COMPILER@") # path +set(HIP_PLATFORM "@HIP_PLATFORM@") #string +set(CMAKE_COMMAND "@CMAKE_COMMAND@") # path +set(HIP_run_make2cmake "@HIP_run_make2cmake@") # path + +@HIP_HOST_FLAGS@ +@_HIP_HIPCC_FLAGS@ +@_HIP_HCC_FLAGS@ +@_HIP_NVCC_FLAGS@ +set(HIP_HIPCC_INCLUDE_ARGS "@HIP_HIPCC_INCLUDE_ARGS@") # list (needs to be in quotes to handle spaces properly) + +set(cmake_dependency_file "@cmake_dependency_file@") # path +set(source_file "@source_file@") # path +set(host_flag "@host_flag@") # bool + +# Determine compiler and compiler flags +if(NOT host_flag) + set(__CC ${HIP_HIPCC_EXECUTABLE}) + if(HIP_PLATFORM STREQUAL "hcc") + set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_HCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_HCC_FLAGS_${build_configuration}}) + else() + set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_NVCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_NVCC_FLAGS_${build_configuration}}) + endif() +else() + set(__CC ${HIP_HOST_COMPILER}) + set(__CC_FLAGS ${CMAKE_HOST_FLAGS} ${CMAKE_HOST_FLAGS_${build_configuration}}) +endif() +set(__CC_INCLUDES ${HIP_HIPCC_INCLUDE_ARGS}) + +# hip_execute_process - Executes a command with optional command echo and status message. +# status - Status message to print if verbose is true +# command - COMMAND argument from the usual execute_process argument structure +# ARGN - Remaining arguments are the command with arguments +# HIP_result - Return value from running the command +macro(hip_execute_process status command) + set(_command ${command}) + if(NOT "x${_command}" STREQUAL "xCOMMAND") + message(FATAL_ERROR "Malformed call to hip_execute_process. Missing COMMAND as second argument. (command = ${command})") + endif() + if(verbose) + execute_process(COMMAND "${CMAKE_COMMAND}" -E echo -- ${status}) + # Build command string to print + set(hip_execute_process_string) + foreach(arg ${ARGN}) + # Escape quotes if any + string(REPLACE "\"" "\\\"" arg ${arg}) + # Surround args with spaces with quotes + if(arg MATCHES " ") + list(APPEND hip_execute_process_string "\"${arg}\"") + else() + list(APPEND hip_execute_process_string ${arg}) + endif() + endforeach() + # Echo the command + execute_process(COMMAND ${CMAKE_COMMAND} -E echo ${hip_execute_process_string}) + endif() + # Run the command + execute_process(COMMAND ${ARGN} RESULT_VARIABLE HIP_result) +endmacro() + +# Delete the target file +hip_execute_process( + "Removing ${generated_file}" + COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}" + ) + +# Generate the dependency file +hip_execute_process( + "Generating dependency file: ${cmake_dependency_file}.pre" + COMMAND "${__CC}" + -M + "${source_file}" + -o "${cmake_dependency_file}.pre" + ${__CC_FLAGS} + ${__CC_INCLUDES} + ) + +if(HIP_result) + message(FATAL_ERROR "Error generating ${generated_file}") +endif() + +# Generate the cmake readable dependency file to a temp file +hip_execute_process( + "Generating temporary cmake readable file: ${cmake_dependency_file}.tmp" + COMMAND "${CMAKE_COMMAND}" + -D "input_file:FILEPATH=${cmake_dependency_file}.pre" + -D "output_file:FILEPATH=${cmake_dependency_file}.tmp" + -D "verbose=${verbose}" + -P "${HIP_run_make2cmake}" + ) + +if(HIP_result) + message(FATAL_ERROR "Error generating ${generated_file}") +endif() + +# Copy the file if it is different +hip_execute_process( + "Copy if different ${cmake_dependency_file}.tmp to ${cmake_dependency_file}" + COMMAND "${CMAKE_COMMAND}" -E copy_if_different "${cmake_dependency_file}.tmp" "${cmake_dependency_file}" + ) + +if(HIP_result) + message(FATAL_ERROR "Error generating ${generated_file}") +endif() + +# Delete the temporary file +hip_execute_process( + "Removing ${cmake_dependency_file}.tmp and ${cmake_dependency_file}.pre" + COMMAND "${CMAKE_COMMAND}" -E remove "${cmake_dependency_file}.tmp" "${cmake_dependency_file}.pre" + ) + +if(HIP_result) + message(FATAL_ERROR "Error generating ${generated_file}") +endif() + +# Generate the output file +hip_execute_process( + "Generating ${generated_file}" + COMMAND "${__CC}" + -c + "${source_file}" + -o "${generated_file}" + ${__CC_FLAGS} + ${__CC_INCLUDES} + ) + +if(HIP_result) + # Make sure that we delete the output file + hip_execute_process( + "Removing ${generated_file}" + COMMAND "${CMAKE_COMMAND}" -E remove "${generated_file}" + ) + message(FATAL_ERROR "Error generating file ${generated_file}") +else() + if(verbose) + message("Generated ${generated_file} successfully.") + endif() +endif() +# vim: ts=4:sw=4:expandtab:smartindent diff --git a/projects/hip/cmake/FindHIP/run_make2cmake.cmake b/projects/hip/cmake/FindHIP/run_make2cmake.cmake new file mode 100644 index 0000000000..d2e3eb5169 --- /dev/null +++ b/projects/hip/cmake/FindHIP/run_make2cmake.cmake @@ -0,0 +1,50 @@ +############################################################################### +# Computes dependencies using HIPCC +############################################################################### + +############################################################################### +# This file converts dependency files generated using hipcc to a format that +# cmake can understand. + +# Input variables: +# +# input_file:STRING=<> Dependency file to parse. Required argument +# output_file:STRING=<> Output file to generate. Required argument + +if(NOT input_file OR NOT output_file) + message(FATAL_ERROR "You must specify input_file and output_file on the command line") +endif() + +file(READ ${input_file} depend_text) + +if (NOT "${depend_text}" STREQUAL "") + string(REPLACE " /" "\n/" depend_text ${depend_text}) + string(REGEX REPLACE "^.*:" "" depend_text ${depend_text}) + string(REGEX REPLACE "[ \\\\]*\n" ";" depend_text ${depend_text}) + + set(dependency_list "") + + foreach(file ${depend_text}) + string(REGEX REPLACE "^ +" "" file ${file}) + if(NOT EXISTS "${file}") + message(WARNING " Removing non-existent dependency file: ${file}") + set(file "") + endif() + + if(NOT IS_DIRECTORY "${file}") + get_filename_component(file_absolute "${file}" ABSOLUTE) + list(APPEND dependency_list "${file_absolute}") + endif() + endforeach() +endif() + +# Remove the duplicate entries and sort them. +list(REMOVE_DUPLICATES dependency_list) +list(SORT dependency_list) + +foreach(file ${dependency_list}) + set(hip_hipcc_depend "${hip_hipcc_depend} \"${file}\"\n") +endforeach() + +file(WRITE ${output_file} "# Generated by: FindHIP.cmake. Do not edit.\nSET(HIP_HIPCC_DEPEND\n ${hip_hipcc_depend})\n\n") +# vim: ts=4:sw=4:expandtab:smartindent diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md index d62f40510a..e21fc984e2 100644 --- a/projects/hip/docs/markdown/hip_faq.md +++ b/projects/hip/docs/markdown/hip_faq.md @@ -213,6 +213,9 @@ export HIP_ATP_MARKER=1 ```shell # Use profile to generate timeline view: /opt/rocm/bin/rocm-profiler -o -A + +Or +/opt/rocm/bin/rocm-profiler -e HIP_ATP_MARKER=1 -o -A ``` #### Using HIP_TRACE_API diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index ba52b88433..628cd36c26 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -669,20 +669,14 @@ The following C++ features are not supported: - Try/catch ## Kernel Compilation -HIP now supports compiling C++/HIP kernels to binary. Eventhough HIP does not support fatbinary (yet), the user can specify the target for which the binary can be generated. The file format for binary is `.co` which means Code Object. The following command builds the binary using `hipcc`. +hipcc now supports compiling C++/HIP kernels to binary code objects. +The user can specify the target for which the binary can be generated. HIP/HCC does not yet support fat binaries so only a single target may be specified. +The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`. -`hipcc --genisa --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` +`hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]` ```[TARGET GPU] = fiji/hawaii [INPUT FILE] = Name of the file containing kernels [OUTPUT FILE] = Name of the generated code object file``` -Note that the kernel file should have `int main(){}` at the end it so that the binary is generated. This happens because HCC generates binaries at linking time rather than compilation - -You need 3 things to run kernel in binary. -1. Kernel Binary -2. Name of kernel binary -3. Name of the kernel - -We already got first two of them. In order to get name of the kernel, try `objdump -x [OUTPUT FILE]`. OUTPUT FILE is file generated by hipcc during kernel compilation. The output from objdump has symbol to the kernel whose name is mangled with `grid_launch_parm`, `__functor`, `__cxxamp_trampoline`. An example of how it looks is `ZN12_GLOBAL__N_137_Z3Cpy16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf` where `Cpy` is the name of the kernel written in C++. diff --git a/projects/hip/docs/markdown/hip_performance.md b/projects/hip/docs/markdown/hip_performance.md index 98197b3db7..bd550c9255 100644 --- a/projects/hip/docs/markdown/hip_performance.md +++ b/projects/hip/docs/markdown/hip_performance.md @@ -12,7 +12,7 @@ There are two possible ways to transfer data from Host to Device (H2D) and Devic #### On Large BAR Setup -There are two possible ways to transfer data from Host to Device (H2D) +There are three possible ways to transfer data from Host to Device (H2D) * Using Staging Buffers * Using PinInPlace * Direct Memcpy @@ -24,12 +24,9 @@ There are two possible ways to transfer data from Host to Device (H2D) Some GPUs may not be able to directly access host memory, and in these cases we need to stage the copy through an optimized pinned staging buffer, to implement H2D and D2H copies.The copy is broken into buffer-sized chunks to limit the size of the buffer and also to provide better performance by overlapping the CPU copies with the DMA copies. -PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA -engine. +PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA engine. -By default staging buffers are used for unpinned memory transfers, however other ways can be used by enabling few environment variables (so no need to build the code again!!!) - -Following environment variables can be used: +By default staging buffers are used for unpinned memory transfers. Environment variables allow control over the unpinned copy algorithm and parameters: - HIP_PININPLACE - This environment variable forces the use of PinInPlace logic for all unpinned memory copies diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 01c36afde4..8b9f1db97b 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -32,7 +32,7 @@ THE SOFTWARE. // #define USE_MEMCPYTOSYMBOL // //Use the new HCC accelerator_view::copy instead of am_copy -#define USE_AV_COPY 0 +#define USE_AV_COPY (__hcc_workweek__ >= 16351) // Compile peer-to-peer support. // >= 2 : use HCC hc:accelerator::get_is_peer @@ -353,18 +353,36 @@ struct LockedBase { }; +class ihipModule_t{ +public: + hsa_executable_t executable; + hsa_code_object_t object; + std::string fileName; + void *ptr; + size_t size; +}; + + +class ihipFunction_t{ +public: + hsa_executable_symbol_t kernel_symbol; + uint64_t kernel; +}; + + template class ihipStreamCriticalBase_t : public LockedBase { public: - ihipStreamCriticalBase_t() : + ihipStreamCriticalBase_t(hc::accelerator_view av) : _last_command_type(ihipCommandCopyH2H), _last_copy_signal(NULL), _signalCursor(0), _oldest_live_sig_id(1), _streamSigId(0), _kernelCnt(0), - _signalCnt(0) + _signalCnt(0), + _av(av) { _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1); }; @@ -395,27 +413,14 @@ public: // 2 are required if a barrier packet is inserted. uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait(). SIGSEQNUM _streamSigId; // Monotonically increasing unique signal id. + + hc::accelerator_view _av; }; typedef ihipStreamCriticalBase_t ihipStreamCritical_t; typedef LockedAccessor LockedAccessor_StreamCrit_t; -class ihipModule_t{ -public: - hsa_executable_t executable; - hsa_code_object_t object; - std::string fileName; - void *ptr; - size_t size; -}; - - -class ihipFunction_t{ -public: - hsa_executable_symbol_t kernel_symbol; - uint64_t kernel; -}; // Internal stream structure. class ihipStream_t { @@ -431,21 +436,31 @@ typedef uint64_t SeqNum_t ; void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); + int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); + //--- - // Thread-safe accessors - these acquire / release mutex: - bool lockopen_preKernelCommand(); + // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex. + LockedAccessor_StreamCrit_t lockopen_preKernelCommand(); void lockclose_postKernelCommand(hc::completion_future &kernel_future); - int preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType); void locked_reclaimSignals(SIGSEQNUM sigNum); void locked_wait(bool assertQueueEmpty=false); - SIGSEQNUM locked_lastCopySeqId() {LockedAccessor_StreamCrit_t crit(_criticalData); return lastCopySeqId(crit); }; + + hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); }; + + void locked_waitEvent(hipEvent_t event); + void locked_recordEvent(hipEvent_t event); + + //--- // Use this if we already have the stream critical data mutex: void wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty=false); - void launchModuleKernel(hsa_signal_t signal, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, uint32_t sharedMemBytes, void *kernarg, size_t kernSize, uint64_t kernel); + void launchModuleKernel(hc::accelerator_view av, hsa_signal_t signal, + uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, + uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, + uint32_t sharedMemBytes, void *kernarg, size_t kernSize, uint64_t kernel); // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function. SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sigId : 0; }; @@ -462,7 +477,6 @@ public: //--- //Public member vars - these are set at initialization and never change: SeqNum_t _id; // monotonic sequence ID - hc::accelerator_view _av; unsigned _flags; @@ -672,7 +686,7 @@ extern const char *ihipErrorString(hipError_t); extern ihipCtx_t *ihipGetTlsDefaultCtx(); extern void ihipSetTlsDefaultCtx(ihipCtx_t *ctx); extern hipError_t ihipSynchronize(void); -extern hipError_t ihipCtxStackUpdate(); +extern void ihipCtxStackUpdate(); extern ihipDevice_t *ihipGetDevice(int); ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex); @@ -704,5 +718,12 @@ inline std::ostream & operator<<(std::ostream& os, const dim3& s) return os; } +// Stream printf functions: +inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e) +{ + os << "event:" << std::hex << static_cast (e); + return os; +} + #endif diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 727604d8d8..f8166e7897 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -50,9 +50,8 @@ THE SOFTWARE. #if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. -#define USE_GRID_LAUNCH_20 1 #else -#define USE_GRID_LAUNCH_20 0 +#error (HCC must support GRID_LAUNCH_20) #endif #define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01) @@ -532,7 +531,7 @@ __device__ void __threadfence_block(void); * * @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. */ -__device__ void __threadfence(void); +__device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional"))); /** * @brief threadfence_system makes writes to pinned system memory visible on host CPU. @@ -543,7 +542,7 @@ __device__ void __threadfence(void); * * @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. */ -__device__ void __threadfence_system(void); +__device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional"))); // doxygen end Memory Fence @@ -633,7 +632,6 @@ extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp); #define KNRM "\x1B[0m" #define KGRN "\x1B[32m" -#if USE_GRID_LAUNCH_20 #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ do {\ grid_launch_parm lp;\ @@ -645,21 +643,6 @@ do {\ _kernelName (lp, ##__VA_ARGS__);\ ihipPostLaunchKernel(trueStream, lp);\ } while(0) -#else -#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ -do {\ - grid_launch_parm lp;\ - lp.groupMemBytes = _groupMemBytes; \ - hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp)); \ - if (HIP_TRACE_API) {\ - fprintf(stderr, KGRN "< inline std::string ToString(hipEvent_t v) { - return ToString(&v); + std::ostringstream ss; + ss << v; + return ss.str(); }; diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index baeb080195..8e9b0d92a7 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -484,6 +484,17 @@ inline static hipError_t hipStreamDestroy(hipStream_t stream) } +inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) +{ + return hipCUDAErrorTohipError(cudaStreamWaitEvent(stream, event, flags)); +} + +inline static hipError_t hipStreamQuery(hipStream_t stream) +{ + return hipCUDAErrorTohipError(cudaStreamQuery(stream)); +} + + inline static hipError_t hipDriverGetVersion(int *driverVersion) { cudaError_t err = cudaDriverGetVersion(driverVersion); @@ -624,6 +635,26 @@ inline static hipError_t hipDeviceGet(hipDevice_t *device, int ordinal) return hipCUResultTohipError(cuDeviceGet(device, ordinal)); } +inline static hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) +{ + return hipCUResultTohipError(cuDeviceComputeCapability(major,minor,device)); +} + +inline static hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device) +{ + return hipCUResultTohipError(cuDeviceGetName(name,len,device)); +} + +inline static hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device) +{ + return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device)); +} + +inline static hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device) +{ + return hipCUResultTohipError(cuDeviceTotalMem(bytes,device)); +} + inline static hipError_t hipModuleLoad(hipModule_t *module, const char* fname) { return hipCUResultTohipError(cuModuleLoad(module, fname)); diff --git a/projects/hip/packaging/hip_base.txt b/projects/hip/packaging/hip_base.txt index 67df3c10f6..e528e688ef 100644 --- a/projects/hip/packaging/hip_base.txt +++ b/projects/hip/packaging/hip_base.txt @@ -4,6 +4,7 @@ project(hip_base) install(DIRECTORY @hip_SOURCE_DIR@/bin DESTINATION . USE_SOURCE_PERMISSIONS) install(DIRECTORY @hip_SOURCE_DIR@/include DESTINATION . PATTERN "hip" EXCLUDE) install(FILES @PROJECT_BINARY_DIR@/.version DESTINATION bin) +install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .) ############################# # Packaging steps diff --git a/projects/hip/samples/0_Intro/hcc_dialects/Makefile b/projects/hip/samples/0_Intro/hcc_dialects/Makefile index fb8fcc0c32..3b5ceca7f0 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/Makefile +++ b/projects/hip/samples/0_Intro/hcc_dialects/Makefile @@ -1,8 +1,9 @@ HCC_HOME?=/opt/rocm/hcc HCC = $(HCC_HOME)/bin/hcc -HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` -HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` +OPT=-O2 +HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT} +HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT} CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I$(HCC_HOME)/include CPPAMP_LDFLAGS= -std=c++amp -L$(HCC_HOME)/lib -Wl,--rpath=$(HCC_HOME)/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive @@ -19,7 +20,7 @@ $(error hcc_dialects requires hcc compiler and only runs on hcc platform) endif -TARGETS=vadd_hc_arrayview vadd_hc_array vadd_amp_arrayview vadd_hip +TARGETS=vadd_hc_arrayview vadd_hc_array vadd_hc_am vadd_amp_arrayview vadd_hip all: $(TARGETS) @@ -51,7 +52,7 @@ vadd_hc_array: vadd_hc_array.o vadd_hc_am.o: vadd_hc_am.cpp $(HCC) $(HCC_CFLAGS) -c $< -o $@ vadd_hc_am: vadd_hc_am.o - $(HCC) $(HCC_LDFLAGS) $< -o $@ + $(HCC) $(HCC_LDFLAGS) -lhc_am $< -o $@ diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp index 5cb2e8c98f..53a137f74c 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp +++ b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp @@ -32,23 +32,25 @@ int main(int argc, char *argv[]) for (int i=0; i (sizeElements), - [&] (hc::index<1> idx) [[hc]] { + [=] (hc::index<1> idx) [[hc]] { int i = idx[0]; C_d[i] = A_d[i] + B_d[i]; }); // This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing. - av.copy(C_d, C_h); // C++ copy D2H + av.copy(C_d, C_h, sizeBytes); // C++ copy D2H for (int i=0; iargBuffer(5); uint32_t *ptr32_t = (uint32_t*)&argBuffer[0]; @@ -79,7 +80,16 @@ int main(){ memcpy(ptr32_t + 5, &one, sizeof(uint32_t)); memcpy(&argBuffer[3], &Ad, sizeof(void*)); memcpy(&argBuffer[4], &Bd, sizeof(void*)); +#endif +#ifdef __HIP_PLATFORM_NVCC__ + uint32_t one = 1; + std::vectorargBuffer(3); + uint32_t *ptr32_t = (uint32_t*)&argBuffer[0]; + memcpy(ptr32_t + 0, &one, sizeof(uint32_t)); + memcpy(&argBuffer[1], &Ad, sizeof(void*)); + memcpy(&argBuffer[2], &Bd, sizeof(void*)); +#endif size_t size = argBuffer.size()*sizeof(void*); diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_isa.cpp b/projects/hip/samples/0_Intro/module_api/vcpy_isa.cpp deleted file mode 100644 index ead3253115..0000000000 --- a/projects/hip/samples/0_Intro/module_api/vcpy_isa.cpp +++ /dev/null @@ -1,9 +0,0 @@ -#include - -__global__ void hello_world(hipLaunchParm lp, float *a, float *b) -{ - int tx = hipThreadIdx_x; - b[tx] = a[tx]; -} - -int main(){} diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_isa.cu b/projects/hip/samples/0_Intro/module_api/vcpy_isa.cu deleted file mode 100644 index d2a0838604..0000000000 --- a/projects/hip/samples/0_Intro/module_api/vcpy_isa.cu +++ /dev/null @@ -1,6 +0,0 @@ - -extern "C" __global__ void hello_world(float *a, float *b) -{ - int tx = threadIdx.x; - b[tx] = a[tx]; -} diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_isa.ptx b/projects/hip/samples/0_Intro/module_api/vcpy_isa.ptx deleted file mode 100644 index 62eb3f63df..0000000000 --- a/projects/hip/samples/0_Intro/module_api/vcpy_isa.ptx +++ /dev/null @@ -1,38 +0,0 @@ -// -// Generated by NVIDIA NVVM Compiler -// -// Compiler Build ID: CL-19856038 -// Cuda compilation tools, release 7.5, V7.5.17 -// Based on LLVM 3.4svn -// - -.version 4.3 -.target sm_20 -.address_size 64 - - // .globl hello_world - -.visible .entry hello_world( - .param .u64 hello_world_param_0, - .param .u64 hello_world_param_1 -) -{ - .reg .f32 %f<2>; - .reg .b32 %r<2>; - .reg .b64 %rd<8>; - - - ld.param.u64 %rd1, [hello_world_param_0]; - ld.param.u64 %rd2, [hello_world_param_1]; - cvta.to.global.u64 %rd3, %rd2; - cvta.to.global.u64 %rd4, %rd1; - mov.u32 %r1, %tid.x; - mul.wide.s32 %rd5, %r1, 4; - add.s64 %rd6, %rd4, %rd5; - ld.global.f32 %f1, [%rd6]; - add.s64 %rd7, %rd3, %rd5; - st.global.f32 [%rd7], %f1; - ret; -} - - diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp new file mode 100644 index 0000000000..0e051f76fc --- /dev/null +++ b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp @@ -0,0 +1,8 @@ +#include "hip/hip_runtime.h" + +extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b) +{ + int tx = hipThreadIdx_x; + b[tx] = a[tx]; +} + diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp index ce829b73d3..bc3677a846 100644 --- a/projects/hip/src/device_util.cpp +++ b/projects/hip/src/device_util.cpp @@ -27,11 +27,6 @@ THE SOFTWARE. using namespace hc::precise_math; #endif -#if __hcc_workweek__ > 16186 -#define USE_DYNAMIC_SHARED 1 -#else -#define USE_DYNAMIC_SHARED 0 -#endif #define HIP_SQRT_2 1.41421356237 #define HIP_SQRT_PI 1.77245385091 @@ -774,11 +769,6 @@ __device__ float __hip_ynf(int n, float x) } -#if __hcc_workweek__ > 16186 -#define USE_DYNAMIC_SHARED 1 -#else -#define USE_DYNAMIC_SHARED 0 -#endif __device__ float acosf(float x) { @@ -1854,12 +1844,10 @@ __host__ __device__ int max(int arg1, int arg2) return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); } -#if USE_DYNAMIC_SHARED __device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr() { return hc::get_dynamic_group_segment_base_pointer(); } -#endif diff --git a/projects/hip/src/hip_context.cpp b/projects/hip/src/hip_context.cpp index e19c45d2c3..dc0a30bd86 100644 --- a/projects/hip/src/hip_context.cpp +++ b/projects/hip/src/hip_context.cpp @@ -29,22 +29,13 @@ THE SOFTWARE. // Stack of contexts thread_local std::stack tls_ctxStack; -hipError_t ihipCtxStackUpdate() +void ihipCtxStackUpdate() { - //HIP_INIT_API(); - hipError_t e = hipSuccess; - if(tls_ctxStack.empty()) { - tls_ctxStack.push(ihipGetTlsDefaultCtx()); + tls_ctxStack.push(ihipGetTlsDefaultCtx()); } - - return ihipLogStatus(e); } -/** - * @return #hipSuccess, #hipErrorInvalidValue - */ -//--- hipError_t hipInit(unsigned int flags) { HIP_INIT_API(flags); @@ -59,10 +50,6 @@ hipError_t hipInit(unsigned int flags) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init @@ -75,10 +62,6 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) return ihipLogStatus(e); } -/** - * @return #hipSuccess, #hipErrorInvalidDevice - */ -//--- hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { HIP_INIT_API(device, deviceId); // FIXME - review if we want to init @@ -93,11 +76,6 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) return ihipLogStatus(e); }; - -/** - * @return #hipSuccess - */ -//--- hipError_t hipDriverGetVersion(int *driverVersion) { HIP_INIT_API(driverVersion); @@ -109,10 +87,6 @@ hipError_t hipDriverGetVersion(int *driverVersion) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess, #hipErrorInvalidValue - */ -//--- hipError_t hipCtxDestroy(hipCtx_t ctx) { HIP_INIT_API(ctx); @@ -135,10 +109,6 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { HIP_INIT_API(ctx); @@ -159,10 +129,6 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess, #hipErrorInvalidContext - */ -//--- hipError_t hipCtxPushCurrent(hipCtx_t ctx) { HIP_INIT_API(ctx); @@ -177,10 +143,6 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { HIP_INIT_API(ctx); @@ -194,10 +156,6 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetCurrent(hipCtx_t ctx) { HIP_INIT_API(ctx); @@ -212,10 +170,6 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess, #hipErrorInvalidContext - */ -//--- hipError_t hipCtxGetDevice(hipDevice_t *device) { HIP_INIT_API(device); @@ -232,10 +186,6 @@ hipError_t hipCtxGetDevice(hipDevice_t *device) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) { HIP_INIT_API(apiVersion); @@ -247,10 +197,6 @@ hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) { HIP_INIT_API(cacheConfig); @@ -260,10 +206,6 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) { HIP_INIT_API(cacheConfig); @@ -273,10 +215,6 @@ hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) { HIP_INIT_API(config); @@ -286,10 +224,6 @@ hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { HIP_INIT_API(pConfig); @@ -299,20 +233,12 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSynchronize ( void ) { HIP_INIT_API(1); return ihipSynchronize(); //TODP Shall check validity of ctx? } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetFlags ( unsigned int* flags ) { HIP_INIT_API(flags); diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 14338c11de..72c92ac76f 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -160,7 +160,7 @@ hipError_t hipSetDevice(int deviceId) */ hipError_t hipDeviceSynchronize(void) { - HIP_INIT_API(1); + HIP_INIT_API(); return ihipSynchronize(); } @@ -324,3 +324,40 @@ hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId) return ihipLogStatus(e); } + +hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) +{ + HIP_INIT_API(major,minor, device); + hipError_t e = hipSuccess; + int deviceId= device->_deviceId; + e = hipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId); + e = hipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId); + return ihipLogStatus(e); +} + +hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device) +{ + HIP_INIT_API(name,len, device); + hipError_t e = hipSuccess; + int nameLen = strlen(device->_props.name); + if(nameLen <= len) + memcpy(name,device->_props.name,nameLen); + return ihipLogStatus(e); +} + +hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device) +{ + HIP_INIT_API(pciBusId,len, device); + hipError_t e = hipSuccess; + int deviceId= device->_deviceId; + e = hipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId); + return ihipLogStatus(e); +} + +hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device) +{ + HIP_INIT_API(bytes, device); + hipError_t e = hipSuccess; + *bytes= device->_props.totalGlobalMem; + return ihipLogStatus(e); +} diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 4324583690..ecc6038b82 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -40,7 +40,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) eh->_flags = flags; eh->_timestamp = 0; eh->_copySeqId = 0; - *event = eh; + *event = eh; // TODO - allocat the event directly, no copy needed. } else { e = hipErrorInvalidValue; } @@ -72,9 +72,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(event, stream); - ihipEvent_t *eh = event; - if (eh && eh->_state != hipEventStatusUnitialized) { - eh->_stream = stream; + if (event && event->_state != hipEventStatusUnitialized) { + event->_stream = stream; if (stream == NULL) { // If stream == NULL, wait on all queues. @@ -83,16 +82,16 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); ctx->locked_syncDefaultStream(true); - eh->_timestamp = hc::get_system_ticks(); - eh->_state = hipEventStatusRecorded; + event->_timestamp = hc::get_system_ticks(); + event->_state = hipEventStatusRecorded; return ihipLogStatus(hipSuccess); } else { - eh->_state = hipEventStatusRecording; + event->_state = hipEventStatusRecording; // Clear timestamps - eh->_timestamp = 0; - eh->_marker = stream->_av.create_marker(); - - eh->_copySeqId = stream->locked_lastCopySeqId(); + event->_timestamp = 0; + + // Record the event in the stream: + stream->locked_recordEvent(event); return ihipLogStatus(hipSuccess); } @@ -122,21 +121,19 @@ hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_API(event); - ihipEvent_t *eh = event; - - if (eh) { - if (eh->_state == hipEventStatusUnitialized) { + if (event) { + if (event->_state == hipEventStatusUnitialized) { return ihipLogStatus(hipErrorInvalidResourceHandle); - } else if (eh->_state == hipEventStatusCreated ) { + } else if (event->_state == hipEventStatusCreated ) { // Created but not actually recorded on any device: return ihipLogStatus(hipSuccess); - } else if (eh->_stream == NULL) { + } else if (event->_stream == NULL) { auto *ctx = ihipGetTlsDefaultCtx(); ctx->locked_syncDefaultStream(true); return ihipLogStatus(hipSuccess); } else { - eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); - eh->_stream->locked_reclaimSignals(eh->_copySeqId); + event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); + event->_stream->locked_reclaimSignals(event->_copySeqId); return ihipLogStatus(hipSuccess); } @@ -196,12 +193,11 @@ hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_API(event); - ihipEvent_t *eh = event; // TODO-stream - need to read state of signal here: The event may have become ready after recording.. // TODO-HCC - use get_hsa_signal here. - if (eh->_state == hipEventStatusRecording) { + if (event->_state == hipEventStatusRecording) { return ihipLogStatus(hipErrorNotReady); } else { return ihipLogStatus(hipSuccess); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index ea8604b7c4..1954b31c70 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -137,7 +137,7 @@ ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex) }; -static thread_local ihipCtx_t *tls_defaultCtx = nullptr; +static thread_local ihipCtx_t *tls_defaultCtx = nullptr; void ihipSetTlsDefaultCtx(ihipCtx_t *ctx) { tls_defaultCtx = ctx; @@ -162,7 +162,7 @@ hipError_t ihipSynchronize(void) { ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. - return ihipLogStatus(hipSuccess); + return (hipSuccess); } //================================================================================================= @@ -195,9 +195,9 @@ ihipSignal_t::~ihipSignal_t() //--- ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags) : _id(0), // will be set by add function. - _av(av), _flags(flags), - _ctx(ctx) + _ctx(ctx), + _criticalData(av) { tprintf(DB_SYNC, " streamCreate: stream=%p\n", this); }; @@ -246,7 +246,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty { if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); - _av.wait(); + crit->_av.wait(); } if (crit->_last_copy_signal) { @@ -273,6 +273,28 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty) }; +#if USE_AV_COPY +// Causes current stream to wait for specified event to complete: +void ihipStream_t::locked_waitEvent(hipEvent_t event) +{ + LockedAccessor_StreamCrit_t crit(_criticalData); + + // TODO - check state of event here: + crit->_av.create_blocking_marker(event->_marker); +} +#endif + +// Create a marker in this stream. +// Save state in the event so it can track the status of the event. +void ihipStream_t::locked_recordEvent(hipEvent_t event) +{ + // Lock the stream to prevent simultaneous access + LockedAccessor_StreamCrit_t crit(_criticalData); + + event->_marker = crit->_av.create_marker(); + event->_copySeqId = lastCopySeqId(crit); +} + //============================================================================= @@ -388,11 +410,10 @@ int HIP_NUM_KERNELS_INFLIGHT = 128; //into the stream to mimic CUDA stream semantics. (some hardware uses separate //queues for data commands and kernel commands, and no implicit ordering is provided). // -bool ihipStream_t::lockopen_preKernelCommand() +LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() { LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/); - bool addedSync = false; if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){ this->wait(crit); @@ -402,9 +423,8 @@ bool ihipStream_t::lockopen_preKernelCommand() // If switching command types, we need to add a barrier packet to synchronize things. if (crit->_last_command_type != ihipCommandKernel) { if (crit->_last_copy_signal) { - addedSync = true; - hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue(); + hsa_queue_t * q = (hsa_queue_t*) (crit->_av.get_hsa_queue()); if (HIP_DISABLE_HW_KERNEL_DEP == 0) { this->enqueueBarrier(q, crit->_last_copy_signal, NULL); tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n", @@ -422,7 +442,7 @@ bool ihipStream_t::lockopen_preKernelCommand() crit->_last_command_type = ihipCommandKernel; } - return addedSync; + return crit; } @@ -433,6 +453,11 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu // We locked _criticalData in the lockopen_preKernelCommand() so OK to access here: _criticalData._last_kernel_future = kernelFuture; + if (HIP_LAUNCH_BLOCKING) { + kernelFuture.wait(); + tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str()); + } + _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand. }; @@ -457,7 +482,7 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t needSync = 1; ihipSignal_t *depSignal = allocSignal(crit); hsa_signal_store_relaxed(depSignal->_hsaSignal,1); - this->enqueueBarrier(static_cast(_av.get_hsa_queue()), NULL, depSignal); + this->enqueueBarrier(static_cast(crit->_av.get_hsa_queue()), NULL, depSignal); *waitSignal = depSignal->_hsaSignal; } else if (crit->_last_copy_signal) { needSync = 1; @@ -487,7 +512,10 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t } -void ihipStream_t::launchModuleKernel(hsa_signal_t signal, +// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av. +void ihipStream_t::launchModuleKernel( + hc::accelerator_view av, + hsa_signal_t signal, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, @@ -500,11 +528,12 @@ void ihipStream_t::launchModuleKernel(hsa_signal_t signal, uint64_t kernel){ hsa_status_t status; void *kern; - hsa_amd_memory_pool_t *pool = reinterpret_cast(_av.get_hsa_kernarg_region()); + + hsa_amd_memory_pool_t *pool = reinterpret_cast(av.get_hsa_kernarg_region()); status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern); - status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)_av.get_hsa_agent(), 0, kern); + status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern); memcpy(kern, kernarg, kernSize); - hsa_queue_t *Queue = (hsa_queue_t*)_av.get_hsa_queue(); + hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue(); const uint32_t queue_mask = Queue->size-1; uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue); hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]); @@ -1117,8 +1146,9 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c while (std::getline(ss, device_id, ',')) { if (atoi(device_id.c_str()) >= 0) { g_hip_visible_devices.push_back(atoi(device_id.c_str())); - }else// Any device number after invalid number will not present + } else { // Any device number after invalid number will not present break; + } } // Print out the number of ids if (HIP_PRINT_ENV) { @@ -1312,11 +1342,11 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) } } -void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream) +void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream) { std::string streamString = ToString(stream); fprintf(stderr, KGRN "<grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z, + kernelName, lp->grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z, lp->dynamic_group_mem_bytes, streamString.c_str());\ } @@ -1327,7 +1357,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ { HIP_INIT(); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; lp->grid_dim.z = grid.z; @@ -1336,27 +1365,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ lp->group_dim.z = block.z; lp->barrier_bit = barrier_bit_queue_default; lp->launch_fence = -1; -#else - lp->gridDim.x = grid.x; - lp->gridDim.y = grid.y; - lp->gridDim.z = grid.z; - lp->groupDim.x = block.x; - lp->groupDim.y = block.y; - lp->groupDim.z = block.z; -#endif - stream->lockopen_preKernelCommand(); -// *av = &stream->_av; - lp->av = &stream->_av; + + auto crit = stream->lockopen_preKernelCommand(); + lp->av = &(crit->_av); lp->cf = new hc::completion_future; -// lp->av = static_cast(av); -// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); } + + hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp) { - HIP_INIT_API(stream, grid, block, lp); + HIP_INIT(); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; lp->grid_dim.y = 1; lp->grid_dim.z = 1; @@ -1365,28 +1385,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri lp->group_dim.z = block.z; lp->barrier_bit = barrier_bit_queue_default; lp->launch_fence = -1; -#else - lp->gridDim.x = grid; - lp->gridDim.y = 1; - lp->gridDim.z = 1; - lp->groupDim.x = block.x; - lp->groupDim.y = block.y; - lp->groupDim.z = block.z; -#endif - stream->lockopen_preKernelCommand(); -// *av = &stream->_av; - lp->av = &stream->_av; + + auto crit = stream->lockopen_preKernelCommand(); + lp->av = &(crit->_av); lp->cf = new hc::completion_future; -// lp->av = static_cast(av); -// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); } + hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp) { - HIP_INIT_API(stream, grid, block, lp); + HIP_INIT(); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; lp->grid_dim.z = grid.z; @@ -1395,28 +1405,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri lp->group_dim.z = 1; lp->barrier_bit = barrier_bit_queue_default; lp->launch_fence = -1; -#else - lp->gridDim.x = grid.x; - lp->gridDim.y = grid.y; - lp->gridDim.z = grid.z; - lp->groupDim.x = block; - lp->groupDim.y = 1; - lp->groupDim.z = 1; -#endif - stream->lockopen_preKernelCommand(); -// *av = &stream->_av; - lp->av = &stream->_av; + + auto crit = stream->lockopen_preKernelCommand(); + lp->av = &(crit->_av); lp->cf = new hc::completion_future; -// lp->av = static_cast(av); -// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); return (stream); } + hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp) { - HIP_INIT_API(stream, grid, block, lp); + HIP_INIT(); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; lp->grid_dim.y = 1; lp->grid_dim.z = 1; @@ -1425,37 +1425,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g lp->group_dim.z = 1; lp->barrier_bit = barrier_bit_queue_default; lp->launch_fence = -1; -#else - lp->gridDim.x = grid; - lp->gridDim.y = 1; - lp->gridDim.z = 1; - lp->groupDim.x = block; - lp->groupDim.y = 1; - lp->groupDim.z = 1; -#endif - stream->lockopen_preKernelCommand(); -// *av = &stream->_av; - lp->av = &stream->_av; - lp->cf = new hc::completion_future; -// lp->av = static_cast(av); -// lp->cf = static_cast(malloc(sizeof(hc::completion_future))); + + auto crit = stream->lockopen_preKernelCommand(); + lp->av = &(crit->_av); + lp->cf = new hc::completion_future; // TODO, is this necessary? return (stream); } //--- //Called after kernel finishes execution. +//This releases the lock on the stream. void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp) { -// stream->lockclose_postKernelCommand(cf); - stream->lockclose_postKernelCommand(*lp.cf); - if (HIP_LAUNCH_BLOCKING) { - tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream); - } + stream->lockclose_postKernelCommand(*(lp.cf)); } -// //================================================================================================= // HIP API Implementation // @@ -1629,7 +1615,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const // TODO - remove, slow path. tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); #if USE_AV_COPY - _av.copy(src,dst,sizeBytes); + crit->_av.copy(src,dst,sizeBytes); #else hc::am_copy(dst, src, sizeBytes); #endif @@ -1677,7 +1663,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const // TODO - remove, slow path. tprintf(DB_COPY1, "D2H && !dstTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); #if USE_AV_COPY - _av.copy(src, dst, sizeBytes); + crit->_av.copy(src, dst, sizeBytes); #else hc::am_copy(dst, src, sizeBytes); #endif @@ -1879,7 +1865,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a stream = device->_defaultStream; } - *av = &(stream->_av); + *av = stream->locked_getAv(); hipError_t err = hipSuccess; return ihipLogStatus(err); diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index a64ad94e81..8516d7520d 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -119,7 +119,13 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) HIP_INIT_API(ptr, sizeBytes); hipError_t hip_status = hipSuccess; - + // return NULL pointer when malloc size is 0 + if (sizeBytes == 0) + { + *ptr = NULL; + return ihipLogStatus(hip_status); + } + auto ctx = ihipGetTlsDefaultCtx(); if (ctx) { @@ -142,6 +148,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hip_status = hipErrorMemoryAllocation; } + //printf (" hipMalloc allocated %p\n", *ptr); + return ihipLogStatus(hip_status); } @@ -672,9 +680,12 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, +// TODO - make member function of stream? template hc::completion_future -ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes) +ihipMemsetKernel(hipStream_t stream, + LockedAccessor_StreamCrit_t &crit, + T * ptr, T val, size_t sizeBytes) { int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits); const int threads_per_wg = 256; @@ -690,7 +701,7 @@ ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes) hc::completion_future cf = hc::parallel_for_each( - stream->_av, + crit->_av, ext_tile, [=] (hc::tiled_index<1> idx) __attribute__((hc)) @@ -707,41 +718,6 @@ ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes) return cf; } -template -hc::completion_future -ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes) -{ - int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits); - const int threads_per_wg = 256; - - int threads = wg * threads_per_wg; - if (threads > sizeBytes) { - threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg; - } - - - hc::extent<1> ext(threads); - auto ext_tile = ext.tile(threads_per_wg); - - hc::completion_future cf = - hc::parallel_for_each( - stream->_av, - ext_tile, - [=] (hc::tiled_index<1> idx) - __attribute__((hc)) - { - int offset = amp_get_global_id(0); - // TODO-HCC - change to hc_get_local_size() - int stride = amp_get_local_size(0) * hc_get_num_groups(0) ; - - for (int i=offset; ilockopen_preKernelCommand(); + auto crit = stream->lockopen_preKernelCommand(); hc::completion_future cf ; @@ -765,7 +741,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s try { value = value & 0xff; unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; - cf = ihipMemsetKernel (stream, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); + cf = ihipMemsetKernel (stream, crit, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); } catch (std::exception &ex) { e = hipErrorInvalidValue; @@ -773,7 +749,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s } else { // use a slow byte-per-workitem copy: try { - cf = ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + cf = ihipMemsetKernel (stream, crit, static_cast (dst), value, sizeBytes); } catch (std::exception &ex) { e = hipErrorInvalidValue; @@ -808,7 +784,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) stream = ihipSyncAndResolveStream(stream); if (stream) { - stream->lockopen_preKernelCommand(); + auto crit = stream->lockopen_preKernelCommand(); hc::completion_future cf ; @@ -817,7 +793,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) try { value = value & 0xff; unsigned value32 = (value << 24) | (value << 16) | (value << 8) | (value) ; - cf = ihipMemsetKernel (stream, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); + cf = ihipMemsetKernel (stream, crit, static_cast (dst), value32, sizeBytes/sizeof(unsigned)); } catch (std::exception &ex) { e = hipErrorInvalidValue; @@ -825,7 +801,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes ) } else { // use a slow byte-per-workitem copy: try { - cf = ihipMemsetKernel (stream, static_cast (dst), value, sizeBytes); + cf = ihipMemsetKernel (stream, crit, static_cast (dst), value, sizeBytes); } catch (std::exception &ex) { e = hipErrorInvalidValue; diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index e43cc62829..28c65b6669 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -104,7 +104,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ *module = new ihipModule_t; if(module == NULL){ - return hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } auto ctx = ihipGetTlsDefaultCtx(); @@ -117,7 +117,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ std::ifstream in(fname, std::ios::binary | std::ios::ate); if(!in){ - return hipErrorFileNotFound; + return ihipLogStatus(hipErrorFileNotFound); }else{ @@ -130,12 +130,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_memory_allocate(sysRegion, size, (void**)&p); if(status != HSA_STATUS_SUCCESS){ - return hipErrorOutOfMemory; + return ihipLogStatus(hipErrorOutOfMemory); } char *ptr = (char*)p; if(!ptr){ - return hipErrorOutOfMemory; + return ihipLogStatus(hipErrorOutOfMemory); } (*module)->ptr = p; (*module)->size = size; @@ -146,27 +146,33 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object); if(status != HSA_STATUS_SUCCESS){ - return hipErrorSharedObjectInitFailed; + return ihipLogStatus(hipErrorSharedObjectInitFailed); } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); if(status != HSA_STATUS_SUCCESS){ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); } } } - return ret; + return ihipLogStatus(ret); } hipError_t hipModuleUnload(hipModule_t hmod){ hipError_t ret = hipSuccess; hsa_status_t status = hsa_executable_destroy(hmod->executable); - if(status != HSA_STATUS_SUCCESS){ret = hipErrorInvalidValue; } + if(status != HSA_STATUS_SUCCESS) + { + ret = hipErrorInvalidValue; + } status = hsa_code_object_destroy(hmod->object); - if(status != HSA_STATUS_SUCCESS){ret = hipErrorInvalidValue; } + if(status != HSA_STATUS_SUCCESS) + { + ret = hipErrorInvalidValue; + } delete hmod; - return ret; + return ihipLogStatus(ret); } hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){ @@ -174,7 +180,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch hipError_t ret = hipSuccess; if(name == nullptr){ - return hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } if(ctx == nullptr){ @@ -189,13 +195,13 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch hsa_status_t status; status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL); if(status != HSA_STATUS_SUCCESS){ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); } status = hsa_executable_freeze(hmod->executable, NULL); status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->kernel_symbol); if(status != HSA_STATUS_SUCCESS){ - return hipErrorNotFound; + return ihipLogStatus(hipErrorNotFound); } status = hsa_executable_symbol_get_info((*func)->kernel_symbol, @@ -203,10 +209,10 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch &(*func)->kernel); if(status != HSA_STATUS_SUCCESS){ - return hipErrorNotFound; + return ihipLogStatus(hipErrorNotFound); } } - return ret; + return ihipLogStatus(ret); } hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, @@ -215,6 +221,7 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, return ihipModuleGetFunction(hfunc, hmod, name); } + hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, @@ -240,10 +247,10 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){ kernSize = *(size_t*)(config[3]); }else{ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); } }else{ - return hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } /* Kernel argument preparation. @@ -262,7 +269,7 @@ Kernel argument preparation. /* Launch AQL packet */ - hStream->launchModuleKernel(signal, blockDimX, blockDimY, blockDimZ, + hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ, gridDimX, gridDimY, gridDimZ, sharedMemBytes, config[1], kernSize, f->kernel); /* @@ -273,10 +280,10 @@ Kernel argument preparation. ihipPostLaunchKernel(hStream, lp); - + } - return ret; + return ihipLogStatus(ret); } @@ -285,17 +292,17 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, HIP_INIT_API(name); hipError_t ret = hipSuccess; if(dptr == NULL || bytes == NULL){ - return hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } if(name == NULL || hmod == NULL){ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); } else{ hipFunction_t func; ihipModuleGetFunction(&func, hmod, name); *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t); *dptr = reinterpret_cast(func->kernel); - return ret; + return ihipLogStatus(ret); } } @@ -303,7 +310,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){ HIP_INIT_API(image); hipError_t ret = hipSuccess; if(image == NULL || module == NULL){ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); }else{ auto ctx = ihipGetTlsDefaultCtx(); *module = new ihipModule_t; @@ -318,12 +325,12 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){ status = hsa_memory_allocate(sysRegion, size, (void**)&p); if(status != HSA_STATUS_SUCCESS){ - return hipErrorOutOfMemory; + return ihipLogStatus(hipErrorOutOfMemory); } char *ptr = (char*)p; if(!ptr){ - return hipErrorOutOfMemory; + return ihipLogStatus(hipErrorOutOfMemory); } (*module)->ptr = p; (*module)->size = size; @@ -333,15 +340,15 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image){ status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object); if(status != HSA_STATUS_SUCCESS){ - return hipErrorSharedObjectInitFailed; + return ihipLogStatus(hipErrorSharedObjectInitFailed); } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); if(status != HSA_STATUS_SUCCESS){ - return hipErrorNotInitialized; + return ihipLogStatus(hipErrorNotInitialized); } } - return ret; + return ihipLogStatus(ret); } diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index a204e2f79a..b5d7d8cf5b 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -66,7 +66,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) } //--- -hipError_t hipStreamCreate(hipStream_t *stream) +hipError_t hipStreamCreate(hipStream_t *stream) { HIP_INIT_API(stream); @@ -74,23 +74,42 @@ hipError_t hipStreamCreate(hipStream_t *stream) } +#if USE_AV_COPY==0 //--- /** * @bug This function conservatively waits for all work in the specified stream to complete. */ +#endif hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags); hipError_t e = hipSuccess; - { - // 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. - stream->locked_wait(); - e = hipSuccess; - } + if (event == nullptr) { + e = hipErrorInvalidResourceHandle; + + } else if (event->_state != hipEventStatusUnitialized) { + + bool fastWait = false; + +#if USE_AV_COPY + if (stream != hipStreamNull) { + stream->locked_waitEvent(event); + + fastWait = true; // don't use the slow host-side synchronization. + } + // TODO - clean up if/else logic when USE_AV_COPY enabled. +#endif + + if (!fastWait) { + // 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. + stream->locked_wait(); + e = hipSuccess; + } + } // else event not recorded, return immediately and don't create marker. return ihipLogStatus(e); }; diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md index c762b09ceb..de73652e43 100644 --- a/projects/hip/tests/README.md +++ b/projects/hip/tests/README.md @@ -50,3 +50,12 @@ ctest ctest -R Memcpy ``` + +### If a test fails: + +Extract the commandline from the testing log: +$ grep -A3 -m2 hipMemcpy-size Testing/Temporary/LastTest.log +36/47 Testing: hipMemcpy-size +36/47 Test: hipMemcpy-size +Command: "/home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory/hipMemcpy" "--tests" "0x6" +Directory: /home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 86ba712ead..eea43659dc 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -166,7 +166,8 @@ build_hip_executable (hipEventRecord hipEventRecord.cpp) build_hip_executable_libcpp (hipHcc hipHcc.cpp) #set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) -build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp) +# __workweek fix. +#build_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp) build_hip_executable (hipHostAlloc hipHostAlloc.cpp) build_hip_executable (hipHostGetFlags hipHostGetFlags.cpp) build_hip_executable (hipHostRegister hipHostRegister.cpp) diff --git a/projects/hip/tests/src/hipModule.cpp b/projects/hip/tests/src/hipModule.cpp index 86ccb3ab9f..44855fb74f 100644 --- a/projects/hip/tests/src/hipModule.cpp +++ b/projects/hip/tests/src/hipModule.cpp @@ -23,6 +23,8 @@ THE SOFTWARE. #include #include +#include "test_common.h" + #define LEN 64 #define SIZE LEN<<2 @@ -45,17 +47,17 @@ int main(){ std::cout< #ifdef __HIP_PLATFORM_HCC__ -//#include "hcc_detail/AM.h" #include "hc_am.hpp" #endif +#define USE_AV_COPY (__hcc_workweek__ >= 16351) + size_t Nbytes = 0; //================================================================================================= @@ -409,11 +410,21 @@ void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir remove if (addDir == Up) { for (char *p = basePtr; p=0; p-=bufferSize) { +#if USE_AV_COPY + hc::AmPointerInfo info(p, p, bufferSize, acc, false, false); + hc::am_memtracker_add(p, info); +#else hc::am_memtracker_add(p, bufferSize, acc, false); +#endif } } diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp index 83ce944189..a7e05c53f4 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp @@ -330,7 +330,7 @@ int main(int argc, char *argv[]) parseMyArguments(argc, argv); - printf ("info: set device to %d\n", p_gpuDevice); + printf ("info: set device to %d tests=%x\n", p_gpuDevice, p_tests); HIPCHECK(hipSetDevice(p_gpuDevice)); if (p_tests & 0x01) { diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp index eeba7fd345..dccc00b0e0 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. #define SIZE 1024*1024*256 int main(){ - float *Ad, *B, *Bd, *Bm, *C, *Cd; + float *Ad, *B, *Bd, *Bm, *C, *Cd, *ptr_0; B = (float*)malloc(SIZE); hipMalloc((void**)&Ad, SIZE); hipHostMalloc((void**)&B, SIZE); @@ -31,12 +31,15 @@ int main(){ hipHostGetDevicePointer((void**)&Cd, C, 0/*flags*/); + HIPCHECK_API(hipMalloc((void**)&ptr_0,0), hipSuccess); + HIPCHECK_API(hipFree(Ad) , hipSuccess); HIPCHECK_API(hipHostFree(Ad) , hipErrorInvalidValue); HIPCHECK_API(hipFree(B) , hipErrorInvalidDevicePointer); // try to hipFree on malloced memory HIPCHECK_API(hipFree(Bd) , hipErrorInvalidDevicePointer); HIPCHECK_API(hipFree(Bm) , hipErrorInvalidDevicePointer); + HIPCHECK_API(hipFree(ptr_0) , hipSuccess); HIPCHECK_API(hipHostFree(Bd) , hipSuccess); HIPCHECK_API(hipHostFree(Bm) , hipSuccess); diff --git a/projects/hip/tests/src/runtimeApi/stream/CMakeLists.txt b/projects/hip/tests/src/runtimeApi/stream/CMakeLists.txt index a40732fbbc..88959d5762 100644 --- a/projects/hip/tests/src/runtimeApi/stream/CMakeLists.txt +++ b/projects/hip/tests/src/runtimeApi/stream/CMakeLists.txt @@ -6,8 +6,10 @@ include_directories( ${HIPTEST_SOURCE_DIR} ) build_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp) build_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp) build_hip_executable (hipStreamL5 hipStreamL5.cpp) +build_hip_executable (hipStreamWaitEvent hipStreamWaitEvent.cpp) # TODO - seg fault #make_test(hipAPIStreamEnable " ") #make_test(hipAPIStreamDisable " ") make_test(hipStreamL5 " ") +make_test(hipStreamWaitEvent " ") diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp new file mode 100644 index 0000000000..c5a74b2bc0 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -0,0 +1,123 @@ +/* +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" +#include "test_common.h" +#include +unsigned p_streams = 6; + + +//------ +// Structure for one stream; +template +class Streamer { +public: + Streamer(size_t numElements); + ~Streamer(); + void runAsync(); + void waitComplete(); + +private: + T *_A_h; + T *_B_h; + T *_C_h; + + T *_A_d; + T *_B_d; + T *_C_d; + + hipStream_t _stream; + hipEvent_t _event; + + size_t _numElements; +}; + +template +Streamer::Streamer(size_t numElements) : + _numElements(numElements) +{ + HipTest::initArrays (&_A_d, &_B_d, &_C_d, &_A_h, &_B_h, &_C_h, numElements, true); + + HIPCHECK(hipStreamCreate(&_stream)); + HIPCHECK(hipEventCreate(&_event)); +}; + +template +void Streamer::runAsync() +{ + printf ("testing: %s numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _B_d, _C_d, _numElements); + HIPCHECK(hipEventRecord(_event, _stream)); + HIPCHECK(hipStreamWaitEvent(_stream, _event, 0)); + +} + + + + +//--- +//Parse arguments specific to this test. +void parseMyArguments(int argc, char *argv[]) +{ + int more_argc = HipTest::parseStandardArguments(argc, argv, false); + + // parse args for this test: + for (int i = 1; i < more_argc; i++) { + const char *arg = argv[i]; + + if (!strcmp(arg, "--streams")) { + if (++i >= argc || !HipTest::parseUInt(argv[i], &p_streams)) { + failed("Bad streams argument"); + } + } else { + failed("Bad argument '%s'", arg); + } + }; +}; + + + +//--- +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + parseMyArguments(argc, argv); + + typedef Streamer FloatStreamer; + + std::vector streamers; + + size_t numElements = N; + + for (int i=0; irunAsync(); + } + + HIPCHECK(hipDeviceSynchronize()); + + passed(); +} diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index cfde590f78..c50bab233b 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -68,8 +68,9 @@ THE SOFTWARE. {\ hipError_t localError = error; \ if (localError != hipSuccess) { \ - printf("%serror: '%s'(%d) at %s:%d%s\n", \ + printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \ KRED,hipGetErrorString(localError), localError,\ + #error,\ __FILE__, __LINE__,KNRM); \ failed("API returned error code.");\ }\ diff --git a/projects/hip/util/vim/hip.vim b/projects/hip/util/vim/hip.vim index c597df6118..17faecfdeb 100644 --- a/projects/hip/util/vim/hip.vim +++ b/projects/hip/util/vim/hip.vim @@ -65,6 +65,7 @@ syn keyword hipFunctionName expf __expf exp logf __logf log syn keyword hipType hipDeviceProp_t syn keyword hipType hipError_t syn keyword hipType hipStream_t +syn keyword hipType hipEvent_t " Runtime functions syn keyword hipFunctionName hipBindTexture hipBindTextureToArray