From d041291293680f01582f808ff029582f0561c615 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 1 Sep 2016 10:39:14 -0500 Subject: [PATCH 01/39] Fixed offline kernel compilation 1. Removed vcpy_isa.ptx as it should be generated during make 2. Made argument padding specific to hcc path 3. Renamed --gencodeobject to --genco 4. Changed Makefile to work on both nvcc and hcc path Change-Id: Ifd053d541085d9ce4fd37bc21b07674786c7163e [ROCm/hip commit: cc8ce4d58cd9470d4c5b5ebce087606d96bbb117] --- projects/hip/bin/hipcc | 4 +- .../hip/samples/0_Intro/module_api/Makefile | 24 +++++++++--- .../samples/0_Intro/module_api/runKernel.cpp | 11 +++++- .../samples/0_Intro/module_api/vcpy_isa.ptx | 38 ------------------- 4 files changed, 29 insertions(+), 48 deletions(-) delete mode 100644 projects/hip/samples/0_Intro/module_api/vcpy_isa.ptx diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 8fde7f426d..422dba4e54 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -181,7 +181,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 +196,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; diff --git a/projects/hip/samples/0_Intro/module_api/Makefile b/projects/hip/samples/0_Intro/module_api/Makefile index 25ae7b8411..db270beaa0 100644 --- a/projects/hip/samples/0_Intro/module_api/Makefile +++ b/projects/hip/samples/0_Intro/module_api/Makefile @@ -6,16 +6,28 @@ HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) ifeq (${HIP_PLATFORM}, hcc) - GENCODEOBJECT_FLAGS=--target-isa-fiji + GENCODEOBJECT_FLAGS=--target-isa=fiji + +vcpy_isa.compile: vcpy_isa.cpp + $(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co + +clean: + rm -f *.co *.out + +endif + +ifeq (${HIP_PLATFORM}, nvcc) + +vcpy_isa.compile: vcpy_isa.cu + $(HIPCC) --genco vcpy_isa.cu -o vcpy_isa.ptx + +clean: + rm -f *.ptx *.out + endif all: vcpy_isa.compile runKernel.hip.out -vcpy_isa.compile: vcpy_isa.cpp - $(HIPCC) --gencodeobject $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co - runKernel.hip.out: runKernel.cpp $(HIPCC) runKernel.cpp -o runKernel.hip.out -clean: - rm -f *.co *.out diff --git a/projects/hip/samples/0_Intro/module_api/runKernel.cpp b/projects/hip/samples/0_Intro/module_api/runKernel.cpp index e4fa1b6d93..dcbaa4cd35 100644 --- a/projects/hip/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/runKernel.cpp @@ -66,8 +66,9 @@ int main(){ hipModuleLoad(&Module, fileName); hipModuleGetFunction(&Function, Module, kernel_name); - uint32_t len = LEN; - uint32_t one = 1; +#ifdef __HIP_PLATFORM_HCC__ + uint32_t len = LEN; + uint32_t one = 1; std::vectorargBuffer(5); uint32_t *ptr32_t = (uint32_t*)&argBuffer[0]; @@ -79,7 +80,13 @@ 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__ + std::vectorargBuffer(2); + memcpy(&argBuffer[0], &Ad, sizeof(void*)); + memcpy(&argBuffer[1], &Bd, sizeof(void*)); +#endif size_t size = argBuffer.size()*sizeof(void*); 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; -} - - From 6d95372587687556ffa395d77b12bc76dd88dcdf Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 1 Sep 2016 18:54:03 +0300 Subject: [PATCH 02/39] [HIPIFY] CUDA Driver API porting to HIP : CUresult enum. enum CUresult was merged with enum cudaError_t into single hipError_t. Thus a majority of HIP error codes has a reflection to Driver's and RT's corresponding error code at the same time. For instance: cuda2hipRename["CUDA_SUCCESS"] = {"hipSuccess", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaSuccess"] = {"hipSuccess", CONV_ERR, API_RUNTIME}; There are a few CUDA return error codes which are RT or Driver specific. For instance: cuda2hipRename["CUDA_ERROR_INVALID_CONTEXT"] = {"hipErrorInvalidContext", CONV_ERR, API_DRIVER}; cuda2hipRename["cudaErrorInvalidMemcpyDirection"] = {"hipErrorInvalidMemcpyDirection", CONV_ERR, API_RUNTIME}; Matchers were changed from "cuda.* | cublas.*" to "cu.*" as CUDA API functions/types starts with 'cu'. [ROCm/hip commit: 43998387aa43e2cb64de59f1a7c9f0ea626a7348] --- projects/hip/clang-hipify/src/Cuda2Hip.cpp | 150 +++++++++++++++++---- 1 file changed, 126 insertions(+), 24 deletions(-) diff --git a/projects/hip/clang-hipify/src/Cuda2Hip.cpp b/projects/hip/clang-hipify/src/Cuda2Hip.cpp index fde5b7e55b..280b1e84cc 100644 --- a/projects/hip/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/hip/clang-hipify/src/Cuda2Hip.cpp @@ -117,30 +117,133 @@ 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}; + // Error API cuda2hipRename["cudaGetLastError"] = {"hipGetLastError", CONV_ERR, API_RUNTIME}; cuda2hipRename["cudaPeekAtLastError"] = {"hipPeekAtLastError", CONV_ERR, API_RUNTIME}; @@ -1038,8 +1141,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 +1811,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 +1822,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 +1830,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"), From 6824b3af6d43261ee741fda32b1476de56739e28 Mon Sep 17 00:00:00 2001 From: pensun Date: Thu, 1 Sep 2016 13:06:55 -0500 Subject: [PATCH 03/39] Modify hipMalloc to take size of 0 and return NULL pointer for it. Modify hipMemoryAllocate test for this change. Change-Id: I884aee0cd5fe0c60aebb4cb37cdbb4833c5bd0c7 [ROCm/hip commit: b726f0e9496a6f6566671c43e1108fe10e28dfae] --- projects/hip/src/hip_memory.cpp | 8 +++++++- .../hip/tests/src/runtimeApi/memory/hipMemoryAllocate.cpp | 5 ++++- 2 files changed, 11 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index a64ad94e81..df1b70a3f8 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) { 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); From aaca975b25be6822280d7b0af7774b9ef33561bd Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 13:28:27 -0500 Subject: [PATCH 04/39] add hipStreamWaitEvent on nvcc [ROCm/hip commit: 32e9436374400ced3feba8cefe447bd11069e2fe] --- projects/hip/include/nvcc_detail/hip_runtime_api.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index baeb080195..59da96a841 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -484,6 +484,12 @@ 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 hipDriverGetVersion(int *driverVersion) { cudaError_t err = cudaDriverGetVersion(driverVersion); From 5703cef9deda30fef74dee477a8fe0038c06ee51 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 2 Sep 2016 12:25:50 +0530 Subject: [PATCH 05/39] Added hipDeviceGet function Change-Id: I7c25bc24680589682af1b199072807523fc37f50 [ROCm/hip commit: d21e58646207ef7c72cda9033790989c793bdcf4] --- projects/hip/include/hcc_detail/hip_runtime_api.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index fbcca3d12f..b3f9340ed6 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -1235,6 +1235,15 @@ hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx); */ hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId); +/** + * @brief Returns a handle to a compute device + * @param [out] device + * @param [in] ordinal + * + * @returns #hipSuccess, #hipErrorInavlidDevice + */ +hipError_t hipDeviceGet(hipDevice_t *device, int ordinal); + /** * @brief Returns the approximate HIP driver version. From bb215fe616506b185ec8cb9bab25e1abf91d9a87 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 2 Sep 2016 14:45:53 +0530 Subject: [PATCH 06/39] Added hipDeviceComputeCapability, hipDeviceGetPCIBusId and hipDeviceGetName Change-Id: Ibe2d975df796712633900ddc7b0734ec2b8ab4ec [ROCm/hip commit: 3f5a2f83972f5d9462835655db08345ca7c1b353] --- .../hip/include/hcc_detail/hip_runtime_api.h | 29 ++++++++++++++++++ .../hip/include/nvcc_detail/hip_runtime_api.h | 15 ++++++++++ projects/hip/src/hip_device.cpp | 30 +++++++++++++++++++ 3 files changed, 74 insertions(+) diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index b3f9340ed6..6bd2e2e982 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -1244,6 +1244,35 @@ hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId); */ hipError_t hipDeviceGet(hipDevice_t *device, int ordinal); +/** + * @brief Returns the compute capability of the device + * @param [out] major + * @param [out] minor + * @param [in] device + * + * @returns #hipSuccess, #hipErrorInavlidDevice + */ +hipError_t hipDeviceComputeCapability(int *major,int *minor,hipDevice_t device); + +/** + * @brief Returns an identifer string for the device. + * @param [out] name + * @param [in] len + * @param [in] device + * + * @returns #hipSuccess, #hipErrorInavlidDevice + */ +hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device); + +/** + * @brief Returns a PCI Bus Id string for the device. + * @param [out] pciBusId + * @param [in] len + * @param [in] device + * + * @returns #hipSuccess, #hipErrorInavlidDevice + */ +hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device); /** * @brief Returns the approximate HIP driver version. diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index 59da96a841..6001226acf 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -630,6 +630,21 @@ 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 hipModuleLoad(hipModule_t *module, const char* fname) { return hipCUResultTohipError(cuModuleLoad(module, fname)); diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 14338c11de..d05f5bc69e 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -324,3 +324,33 @@ 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); +} + From df654fa6232091fa2a7d47bd56c0fd5d88dabb2e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 2 Sep 2016 15:07:33 +0530 Subject: [PATCH 07/39] FindHIP: First version of re-implementation Change-Id: I84332a0f7d89271bf9370dbcd19d893119e738f9 [ROCm/hip commit: b0621a4794925682a62482a90f7cc34587b4a658] --- projects/hip/bin/hipcc | 9 + projects/hip/cmake/CMakeLists.txt | 12 - projects/hip/cmake/FindHCC.cmake | 73 --- projects/hip/cmake/FindHIP.cmake | 509 ++++++++++++++++-- projects/hip/cmake/FindHIP/run_hipcc.cmake | 163 ++++++ .../hip/cmake/FindHIP/run_make2cmake.cmake | 50 ++ 6 files changed, 680 insertions(+), 136 deletions(-) delete mode 100644 projects/hip/cmake/CMakeLists.txt delete mode 100644 projects/hip/cmake/FindHCC.cmake create mode 100644 projects/hip/cmake/FindHIP/run_hipcc.cmake create mode 100644 projects/hip/cmake/FindHIP/run_make2cmake.cmake diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 422dba4e54..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 = (); @@ -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/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 From 6253a470535c8cd67b455031ebca6efcecf285bb Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 2 Sep 2016 15:08:36 +0530 Subject: [PATCH 08/39] CMakeLists.txt: install cmake folder Change-Id: I5367281c62529b4b832c77ea93e3994a5f9133f4 [ROCm/hip commit: 295df9e006af019ce83a078242cbef8172dff365] --- projects/hip/CMakeLists.txt | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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() ############################# From bb888d709602ea452fe8dd4d2aa3631e533c489a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 2 Sep 2016 15:09:05 +0530 Subject: [PATCH 09/39] hip_base package: bundle cmake folder in package Change-Id: I7f30469242a0357c68b9a27b7b18bbe1f1f15fc1 [ROCm/hip commit: bc055d960a1e3b4ae9d8dbe1be16b633cbf5bd89] --- projects/hip/packaging/hip_base.txt | 1 + 1 file changed, 1 insertion(+) 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 From af9fcb7c87a6dc5719e70058d3f9d1e297d192c1 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 2 Sep 2016 15:11:36 +0530 Subject: [PATCH 10/39] Added hipDeviceTotalMem Change-Id: I877fbc9e4767bbd70ecd0184f5123c9bc6cbd06f [ROCm/hip commit: 2f8d9e20b2d11353a4eedade63c0ac454dfda75d] --- projects/hip/include/hcc_detail/hip_runtime_api.h | 9 +++++++++ projects/hip/include/nvcc_detail/hip_runtime_api.h | 5 +++++ projects/hip/src/hip_device.cpp | 7 +++++++ 3 files changed, 21 insertions(+) diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 6bd2e2e982..029682a341 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -1274,6 +1274,15 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device); */ hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device); +/** + * @brief Returns the total amount of memory on the device. + * @param [out] bytes + * @param [in] device + * + * @returns #hipSuccess, #hipErrorInavlidDevice + */ +hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device); + /** * @brief Returns the approximate HIP driver version. * diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index 6001226acf..e2bcf41dd9 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -645,6 +645,11 @@ inline static hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t 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/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index d05f5bc69e..c3acbadbff 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -354,3 +354,10 @@ hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device) 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); +} From 2d6f4cfb81472ba5207928fa3c3d0745e8100184 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 2 Sep 2016 18:37:26 +0530 Subject: [PATCH 11/39] Removed Doxygen return codes from source file Change-Id: Iaf9f13df73a9c17c494c9ff67203b836ee90648d [ROCm/hip commit: eada9b4444d98b656184ea0e9e029d86d4936893] --- projects/hip/src/hip_context.cpp | 69 -------------------------------- 1 file changed, 69 deletions(-) diff --git a/projects/hip/src/hip_context.cpp b/projects/hip/src/hip_context.cpp index e19c45d2c3..4bf28ad191 100644 --- a/projects/hip/src/hip_context.cpp +++ b/projects/hip/src/hip_context.cpp @@ -41,10 +41,6 @@ hipError_t ihipCtxStackUpdate() return ihipLogStatus(e); } -/** - * @return #hipSuccess, #hipErrorInvalidValue - */ -//--- hipError_t hipInit(unsigned int flags) { HIP_INIT_API(flags); @@ -59,10 +55,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 +67,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 +81,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 +92,6 @@ hipError_t hipDriverGetVersion(int *driverVersion) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess, #hipErrorInvalidValue - */ -//--- hipError_t hipCtxDestroy(hipCtx_t ctx) { HIP_INIT_API(ctx); @@ -135,10 +114,6 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { HIP_INIT_API(ctx); @@ -159,10 +134,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 +148,6 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { HIP_INIT_API(ctx); @@ -194,10 +161,6 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx) return ihipLogStatus(e); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetCurrent(hipCtx_t ctx) { HIP_INIT_API(ctx); @@ -212,10 +175,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 +191,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 +202,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 +211,6 @@ hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) { HIP_INIT_API(cacheConfig); @@ -273,10 +220,6 @@ hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) { HIP_INIT_API(config); @@ -286,10 +229,6 @@ hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) return ihipLogStatus(hipSuccess); } -/** - * @return #hipSuccess - */ -//--- hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) { HIP_INIT_API(pConfig); @@ -299,20 +238,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); From cac6e516ac4dbde4739de9f6570456df38c39415 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 2 Sep 2016 16:31:46 +0300 Subject: [PATCH 12/39] [HIPIFY] CUDA Driver API porting to HIP : types are added. CUdevice -> hipDevice_t // Yet under discussion CUdevice_attribute -> hipDeviceAttribute_t CUdevprop -> hipDeviceProp_t CUfunction -> hipFunction_t CUfunc_cache -> hipFuncCache CUsharedconfig -> hipSharedMemConfig CUcontext -> hipCtx_t CUmodule -> hipModule_t CUevent -> hipEvent_t CUstream -> hipStream_t [ROCm/hip commit: f7818c4c440b680744f673a37a421e99b20101ed] --- projects/hip/clang-hipify/src/Cuda2Hip.cpp | 34 ++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/projects/hip/clang-hipify/src/Cuda2Hip.cpp b/projects/hip/clang-hipify/src/Cuda2Hip.cpp index 280b1e84cc..474f1a50dc 100644 --- a/projects/hip/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/hip/clang-hipify/src/Cuda2Hip.cpp @@ -244,6 +244,40 @@ struct cuda2hipMap { 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["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["CUsharedconfig_enum"] = {"hipSharedMemConfig", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUsharedconfig"] = {"hipSharedMemConfig", CONV_TYPE, 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}; From 3f22169885dc7419216f2cf435582d26760e5d69 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 2 Sep 2016 09:15:49 -0500 Subject: [PATCH 13/39] Documentation review and update Change-Id: If40d096646014d70a198db3532758028abe6a93f [ROCm/hip commit: b1b099941d6361daa3da117121ec8ce5a370574f] --- projects/hip/CONTRIBUTING.md | 3 ++- .../hip/docs/markdown/hip_kernel_language.md | 21 +++++++++++-------- projects/hip/docs/markdown/hip_performance.md | 9 +++----- .../hip/include/hcc_detail/hip_runtime_api.h | 5 ++--- 4 files changed, 19 insertions(+), 19 deletions(-) diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index 060e70e519..6adf73bdf5 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. diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index ba52b88433..c9df5ce1b1 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -669,20 +669,23 @@ 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 +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++. +To load a kernel into HIP, we need both the code object and the name of the kernel stored within the code object. +In order to get name of the kernel, use: +``` +$ objdump -x [CODE_OBJECT_FILE]`. +``` +CODE_OBJECT_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++. The hipLoadKernelModule API needs to specify this mangled name on the HIP/hcc path. + 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_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 029682a341..213928151a 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -726,7 +726,7 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags) ; /** - * @brief Get flags associated with host pointer + * @brief Return flags associated with host pointer * * @param[out] flagsPtr Memory location to store flags * @param[in] hostPtr Host Pointer allocated through hipHostMalloc @@ -1186,13 +1186,12 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ); hipError_t hipCtxSynchronize ( void ); /** - * @brief Get flags used for creating current/default context. + * @brief Return flags used for creating default context. * * @param [out] flags * * @returns #hipSuccess. */ - hipError_t hipCtxGetFlags ( unsigned int* flags ); /** From d92b14aa219175ff0d835a61a14c551be941c0f3 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 2 Sep 2016 09:31:37 -0500 Subject: [PATCH 14/39] fixed code according to convention Change-Id: Ib90c84a30ea6da3531e8b7c3e281dd4b31915198 [ROCm/hip commit: 7a807d54dfac332ca068d721a084d6923af0518a] --- projects/hip/src/hip_module.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index e43cc62829..f81c5efa67 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -162,9 +162,15 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ 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; } From 3e26dd5b5d62409adbd6187b7667a047591c9217 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 2 Sep 2016 09:44:00 -0500 Subject: [PATCH 15/39] added logger for return status Change-Id: Ib93012a15b1cfa90714a6ebd5153cc40e28e6566 [ROCm/hip commit: 04f97aa956c4c9a3cadc291e287857876f8b6285] --- projects/hip/src/hip_module.cpp | 50 ++++++++++++++++----------------- 1 file changed, 25 insertions(+), 25 deletions(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index f81c5efa67..456d3725b8 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,17 +146,17 @@ 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){ @@ -172,7 +172,7 @@ hipError_t hipModuleUnload(hipModule_t hmod){ ret = hipErrorInvalidValue; } delete hmod; - return ret; + return ihipLogStatus(ret); } hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){ @@ -180,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){ @@ -195,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, @@ -209,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, @@ -246,10 +246,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. @@ -282,7 +282,7 @@ Kernel argument preparation. } - return ret; + return ihipLogStatus(ret); } @@ -291,17 +291,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); } } @@ -309,7 +309,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; @@ -324,12 +324,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; @@ -339,15 +339,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); } From 24917dd6e812d9a5c311c7889e8f155addce23c4 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 29 Aug 2016 21:59:37 -0500 Subject: [PATCH 16/39] enable dynamic shared always Change-Id: Idc246546f583a82489ef34d1fe22d0e118b1e0ed [ROCm/hip commit: bbbcb4efad7c39264ac8638c7e1be257b4097780] --- projects/hip/src/device_util.cpp | 12 ------------ 1 file changed, 12 deletions(-) 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 From 1ef87ac5f2dcb4e9983bb01993ae190b60cda2f3 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 29 Aug 2016 22:03:00 -0500 Subject: [PATCH 17/39] enable hc_am example in hcc_ddialects example Change-Id: Iec2f9eb05f95cb025c157fee8fd284aab844d1a2 [ROCm/hip commit: cd3a0a2d61a5e96d298dd9b93ec2005633fbc59f] --- .../hip/samples/0_Intro/hcc_dialects/Makefile | 9 +++++---- .../samples/0_Intro/hcc_dialects/vadd_hc_am.cpp | 15 +++++++++++---- 2 files changed, 16 insertions(+), 8 deletions(-) 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..3685e688cd 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,30 @@ int main(int argc, char *argv[]) for (int i=0; i (sizeElements), [&] (hc::index<1> idx) [[hc]] { int i = idx[0]; - C_d[i] = A_d[i] + B_d[i]; + //C_d[i] = A_d[i] + B_d[i]; + C_d[0] = A_d[1] + B_d[2]; }); + cf.wait(); +#endif + // 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; i Date: Tue, 30 Aug 2016 17:29:50 -0500 Subject: [PATCH 18/39] Refactor for stream->_av. - move _av into stream critical section. ( HCC accelerator_view is not thread-safe but HIP steram is. ) - Refactored many places in code that need to acquire critical section. some were previously thread races, ie enqueueing marker. -remove support for GRID_LAUNCH_VERSION < 20 -Enable USE_AV_COPY based on HCC work-week. - Review hipModule docs, some calrity/editing. Change-Id: I3ce7c25ece048c3504f55ecd4683e506bb1fc8b6 [ROCm/hip commit: e76a272d48138acdea06bb98089b57e4df214787] --- projects/hip/include/hcc_detail/hip_hcc.h | 57 ++++--- projects/hip/include/hcc_detail/hip_runtime.h | 19 +-- .../hip/include/hcc_detail/hip_runtime_api.h | 16 +- projects/hip/src/hip_event.cpp | 19 ++- projects/hip/src/hip_hcc.cpp | 142 ++++++++---------- projects/hip/src/hip_memory.cpp | 54 ++----- projects/hip/src/hip_stream.cpp | 4 + 7 files changed, 128 insertions(+), 183 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 01c36afde4..fb10af1106 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,16 +436,23 @@ 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); @@ -462,7 +474,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; diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 727604d8d8..108f73a7eb 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) @@ -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 "<_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); } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index ea8604b7c4..ca650e5283 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -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,18 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty) }; + +// 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 +400,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 +413,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 +432,7 @@ bool ihipStream_t::lockopen_preKernelCommand() crit->_last_command_type = ihipCommandKernel; } - return addedSync; + return crit; } @@ -433,6 +443,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 +472,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; @@ -500,11 +515,17 @@ 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()); + + // Lock stream to prevent other threads from enqueueing kernels at same time. + LockedAccessor_StreamCrit_t crit (_criticalData); + + hc::accelerator_view av = crit->_av; + + 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 +1138,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) { @@ -1327,7 +1349,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 +1357,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 +1377,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 +1397,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 +1417,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 +1607,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 +1655,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 +1857,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 df1b70a3f8..beb84bdb6f 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -678,9 +678,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; @@ -696,7 +699,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)) @@ -713,41 +716,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 ; @@ -771,7 +739,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; @@ -779,7 +747,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; @@ -814,7 +782,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 ; @@ -823,7 +791,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; @@ -831,7 +799,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_stream.cpp b/projects/hip/src/hip_stream.cpp index a204e2f79a..5d900fcaeb 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -84,6 +84,10 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int hipError_t e = hipSuccess; +#ifdef USE_AV_COPY + printf ("USE_AV_COPY\n"); +#endif + { // 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. From bcd1bdb1328b356673a587fbe78762af12ee96e5 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 13:58:54 -0500 Subject: [PATCH 19/39] improve debug for event Change-Id: If6c0e25729d9964cbe5967833247a6d791d575a0 [ROCm/hip commit: d1fb90966e4f9d8627c6b77712f5cb30b099bffd] --- projects/hip/include/hcc_detail/hip_hcc.h | 7 +++++++ projects/hip/include/hcc_detail/trace_helper.h | 4 +++- 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index fb10af1106..216e81c67f 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -715,5 +715,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/trace_helper.h b/projects/hip/include/hcc_detail/trace_helper.h index 3740d8b9a7..7a8f1106f0 100644 --- a/projects/hip/include/hcc_detail/trace_helper.h +++ b/projects/hip/include/hcc_detail/trace_helper.h @@ -63,7 +63,9 @@ inline std::string ToString(T v) template <> inline std::string ToString(hipEvent_t v) { - return ToString(&v); + std::ostringstream ss; + ss << v; + return ss.str(); }; From 71f898d534d3571b75579865776f776c9fd51b1a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 13:59:55 -0500 Subject: [PATCH 20/39] Use create_blocking_marker for WaitEvent implementation Change-Id: Ib3113f69a14e48b9fe0558d7b455148e478d8eed [ROCm/hip commit: aa823871dbc4062f490596f03c10d5479cb01c70] --- projects/hip/src/hip_event.cpp | 19 ++++++++----------- projects/hip/src/hip_hcc.cpp | 10 ++++++++++ projects/hip/src/hip_stream.cpp | 32 ++++++++++++++++++++++++-------- 3 files changed, 42 insertions(+), 19 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 33ec0bc897..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; } @@ -121,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); } @@ -195,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 ca650e5283..1854e3d8c7 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -273,6 +273,16 @@ 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. diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index 5d900fcaeb..f400e27799 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -74,27 +74,43 @@ hipError_t hipStreamCreate(hipStream_t *stream) } +#ifndef USE_AV_COPY //--- /** * @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; + if (event == nullptr) { + e = hipErrorInvalidResourceHandle; + + } else if (event->_state != hipEventStatusUnitialized) { + + bool fastWait = false; + #ifdef USE_AV_COPY - printf ("USE_AV_COPY\n"); + if (stream != hipStreamNull) { + printf ("HIP: wait locked stream\n"); + 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 - { - // 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 (!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); }; From 4eb8d5c171dd0217059ce1c4ab680f11aed9e5c2 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 14:00:46 -0500 Subject: [PATCH 21/39] Fix av::copy in dialects to use capture-by-value Change-Id: Ibce1488a1326f66b92b4d5b351230666b691ed31 [ROCm/hip commit: fb7046160fe059d25307b0b2fbb047eeb38b6fb1] --- projects/hip/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp | 9 ++------- projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp | 7 +++++-- 2 files changed, 7 insertions(+), 9 deletions(-) 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 3685e688cd..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 @@ -41,17 +41,12 @@ int main(int argc, char *argv[]) // Launch kernel onto AV. // Because the kernel PFE and the copies are submitted to same AV, they will execute in order // and we don't need additional synchronization to ensure the copies complete before the PFE begins. -#if 1 hc::completion_future cf= hc::parallel_for_each(av, hc::extent<1> (sizeElements), - [&] (hc::index<1> idx) [[hc]] { + [=] (hc::index<1> idx) [[hc]] { int i = idx[0]; - //C_d[i] = A_d[i] + B_d[i]; - C_d[0] = A_d[1] + B_d[2]; + C_d[i] = A_d[i] + B_d[i]; }); - cf.wait(); -#endif - // This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing. diff --git a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp index 9d223ba276..9684bef140 100644 --- a/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp +++ b/projects/hip/samples/0_Intro/hcc_dialects/vadd_hip.cpp @@ -27,22 +27,25 @@ int main(int argc, char *argv[]) hipMalloc(&B_d, sizeBytes); hipMalloc(&C_d, sizeBytes); - // Initialize host data + // Initialize host memory for (int i=0; i Date: Thu, 1 Sep 2016 15:43:44 -0500 Subject: [PATCH 22/39] Add hipEvent_t to vim syn Change-Id: Iade4aa123e72ada39028543890957cec03dd7a2e [ROCm/hip commit: 7036b7047b31cdf768496fda4b797b07bef5def1] --- projects/hip/util/vim/hip.vim | 1 + 1 file changed, 1 insertion(+) 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 From 85ae05c401f52936b696bfc42e0a56904e35035e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 15:44:39 -0500 Subject: [PATCH 23/39] Add test for hipStreamWaitEvent Change-Id: I157b33a78a0fc125b4b85f611f32f3b492f1b4d9 [ROCm/hip commit: 472999f5a3bb2277f9c5cc5db431da783fd87f38] --- .../src/runtimeApi/stream/CMakeLists.txt | 2 + .../runtimeApi/stream/hipStreamWaitEvent.cpp | 120 ++++++++++++++++++ 2 files changed, 122 insertions(+) create mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp 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..63685bb509 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -0,0 +1,120 @@ +/* +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" +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\n", __func__, _numElements); + 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()); +} From d3fb92374d1cee210e63e9c8d29de241c209bb63 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 17:58:24 -0500 Subject: [PATCH 24/39] Only use ihipLogStatus from top-level HIP functions. Change-Id: I07e9c088d5c16a79ed52cb008a798889a656016c [ROCm/hip commit: db9fe9f49428cae226ae79ce34142ffe459f47ad] --- projects/hip/CONTRIBUTING.md | 2 ++ projects/hip/include/hcc_detail/hip_hcc.h | 2 +- projects/hip/src/hip_context.cpp | 9 ++------- projects/hip/src/hip_device.cpp | 2 +- projects/hip/src/hip_hcc.cpp | 2 +- 5 files changed, 7 insertions(+), 10 deletions(-) diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index 6adf73bdf5..7858ef58eb 100644 --- a/projects/hip/CONTRIBUTING.md +++ b/projects/hip/CONTRIBUTING.md @@ -121,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/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 216e81c67f..3ad1efddeb 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -683,7 +683,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); diff --git a/projects/hip/src/hip_context.cpp b/projects/hip/src/hip_context.cpp index 4bf28ad191..dc0a30bd86 100644 --- a/projects/hip/src/hip_context.cpp +++ b/projects/hip/src/hip_context.cpp @@ -29,16 +29,11 @@ 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); } hipError_t hipInit(unsigned int flags) diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index c3acbadbff..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(); } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 1854e3d8c7..cb47dafc7d 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -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); } //================================================================================================= From 1edcb0670129c822cf4fecf7ff5bd80225ef19df Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 1 Sep 2016 18:00:31 -0500 Subject: [PATCH 25/39] hipStreamWaitEvent update. Add passed to hipStreamWaitEvent test. Fix pointerAttrib test to work with new and old HCC versions. Minor code touchup. Change-Id: I139ba6ce9f6bf2b4bee89aebdec5981b4346ffc0 [ROCm/hip commit: 0923c2d261f2306d2489a8afed1f2c57b7f47be4] --- projects/hip/src/hip_memory.cpp | 2 ++ projects/hip/src/hip_stream.cpp | 4 ++-- projects/hip/tests/src/hipPointerAttrib.cpp | 13 ++++++++++++- .../src/runtimeApi/stream/hipStreamWaitEvent.cpp | 4 +++- 4 files changed, 19 insertions(+), 4 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index beb84bdb6f..8516d7520d 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -148,6 +148,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hip_status = hipErrorMemoryAllocation; } + //printf (" hipMalloc allocated %p\n", *ptr); + return ihipLogStatus(hip_status); } diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index f400e27799..e044b6b578 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -74,7 +74,7 @@ hipError_t hipStreamCreate(hipStream_t *stream) } -#ifndef USE_AV_COPY +#if USE_AV_COPY==0 //--- /** * @bug This function conservatively waits for all work in the specified stream to complete. @@ -93,7 +93,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int bool fastWait = false; -#ifdef USE_AV_COPY +#if USE_AV_COPY if (stream != hipStreamNull) { printf ("HIP: wait locked stream\n"); stream->locked_waitEvent(event); diff --git a/projects/hip/tests/src/hipPointerAttrib.cpp b/projects/hip/tests/src/hipPointerAttrib.cpp index ab620eb022..f0d04f81bf 100644 --- a/projects/hip/tests/src/hipPointerAttrib.cpp +++ b/projects/hip/tests/src/hipPointerAttrib.cpp @@ -28,11 +28,12 @@ THE SOFTWARE. #include #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/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 63685bb509..7bdae6a96f 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -62,7 +62,7 @@ Streamer::Streamer(size_t numElements) : template void Streamer::runAsync() { - printf ("testing: %s numElements=%zu\n", __func__, _numElements); + 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)); @@ -117,4 +117,6 @@ int main(int argc, char *argv[]) } HIPCHECK(hipDeviceSynchronize()); + + passed(); } From 1b406221432959a787bff6810de1b4655795a887 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 2 Sep 2016 09:38:42 -0500 Subject: [PATCH 26/39] Doc and test updates. Change-Id: I44cbd2c0de7cb692d0194d1bb8b8ae25c58d4e00 [ROCm/hip commit: 84b14aa9d62c29bc63a14940a656291c8c143967] --- projects/hip/docs/markdown/hip_faq.md | 3 +++ projects/hip/tests/README.md | 9 ++++++++ projects/hip/tests/src/CMakeLists.txt | 3 ++- projects/hip/tests/src/hipModule.cpp | 21 +++++++++++-------- .../src/runtimeApi/memory/hipMemcpyAsync.cpp | 2 +- projects/hip/tests/src/test_common.h | 3 ++- 6 files changed, 29 insertions(+), 12 deletions(-) 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/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< Date: Fri, 2 Sep 2016 09:50:49 -0500 Subject: [PATCH 27/39] Add deprecated information for threadfence() and threadfence_system() function. Change-Id: I71793e9f48523776a87c91f3ed1eef7b9a9da88b [ROCm/hip commit: 1f2607153f92ac3a9ad9841c34b16e7dd15c21cc] --- projects/hip/include/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 108f73a7eb..f8166e7897 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -531,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. @@ -542,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 From 58206709107ad25ed582f55686b0233e223a870a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 2 Sep 2016 12:47:25 -0500 Subject: [PATCH 28/39] Fix double-lock of stream on hipModuleLaunchKernel Change-Id: I4ca164971c25f4eb8fbcca11d6258367bb3d2ab4 [ROCm/hip commit: cdba60a56619fca0a14ca823238c13ef2eb98154] --- projects/hip/include/hcc_detail/hip_hcc.h | 7 +++++-- projects/hip/samples/0_Intro/module_api/Makefile | 5 ++++- projects/hip/src/hip_hcc.cpp | 16 +++++++--------- projects/hip/src/hip_module.cpp | 5 +++-- 4 files changed, 19 insertions(+), 14 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 3ad1efddeb..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 (__hcc_workweek__ >= 16351) +#define USE_AV_COPY (__hcc_workweek__ >= 16351) // Compile peer-to-peer support. // >= 2 : use HCC hc:accelerator::get_is_peer @@ -457,7 +457,10 @@ typedef uint64_t SeqNum_t ; // 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; }; diff --git a/projects/hip/samples/0_Intro/module_api/Makefile b/projects/hip/samples/0_Intro/module_api/Makefile index db270beaa0..8981938e91 100644 --- a/projects/hip/samples/0_Intro/module_api/Makefile +++ b/projects/hip/samples/0_Intro/module_api/Makefile @@ -3,11 +3,14 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc +OPT= HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) ifeq (${HIP_PLATFORM}, hcc) GENCODEOBJECT_FLAGS=--target-isa=fiji +all: runKernel.hip.out + vcpy_isa.compile: vcpy_isa.cpp $(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co @@ -29,5 +32,5 @@ endif all: vcpy_isa.compile runKernel.hip.out runKernel.hip.out: runKernel.cpp - $(HIPCC) runKernel.cpp -o runKernel.hip.out + $(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index cb47dafc7d..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; @@ -512,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, @@ -526,11 +529,6 @@ void ihipStream_t::launchModuleKernel(hsa_signal_t signal, hsa_status_t status; void *kern; - // Lock stream to prevent other threads from enqueueing kernels at same time. - LockedAccessor_StreamCrit_t crit (_criticalData); - - hc::accelerator_view av = crit->_av; - 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); @@ -1344,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());\ } diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 456d3725b8..28c65b6669 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -221,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, @@ -268,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); /* @@ -279,7 +280,7 @@ Kernel argument preparation. ihipPostLaunchKernel(hStream, lp); - + } return ihipLogStatus(ret); From 6e9fc7523d813c283d6f7f705d43e0850498467c Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 2 Sep 2016 13:17:17 -0500 Subject: [PATCH 29/39] corrected offline kernel compilation on hipcc path 1. hipgenisa.sh now adds int main(){} during kernel compilation. User does not have to put it there 2. Renamed vcpy_isa.cpp to vcpy_kernel.cpp 3. Removed vcpy_isa.cu as the kernel code should be common for both paths 4. Changed Makefile and runkernel.cpp to work with above changes Change-Id: I9f8c84706b44bb500bc493a68e959762b55a0142 [ROCm/hip commit: 4f8c45d05a552665436c9866d39277b33d469135] --- projects/hip/bin/hipgenisa.sh | 7 ++++++- projects/hip/samples/0_Intro/module_api/Makefile | 10 +++++----- projects/hip/samples/0_Intro/module_api/runKernel.cpp | 8 ++++---- projects/hip/samples/0_Intro/module_api/vcpy_isa.cu | 6 ------ .../module_api/{vcpy_isa.cpp => vcpy_kernel.cpp} | 1 - 5 files changed, 15 insertions(+), 17 deletions(-) delete mode 100644 projects/hip/samples/0_Intro/module_api/vcpy_isa.cu rename projects/hip/samples/0_Intro/module_api/{vcpy_isa.cpp => vcpy_kernel.cpp} (91%) diff --git a/projects/hip/bin/hipgenisa.sh b/projects/hip/bin/hipgenisa.sh index e60abbf78f..03952f49d8 100755 --- a/projects/hip/bin/hipgenisa.sh +++ b/projects/hip/bin/hipgenisa.sh @@ -22,10 +22,15 @@ 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 +cp $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 /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 $FILE_NAMES.kernel.tmp.cpp rm -r /tmp/hipgenisa export KMDUMPISA=0 export KMDUMPLLVM=0 diff --git a/projects/hip/samples/0_Intro/module_api/Makefile b/projects/hip/samples/0_Intro/module_api/Makefile index 8981938e91..99cff6bc0a 100644 --- a/projects/hip/samples/0_Intro/module_api/Makefile +++ b/projects/hip/samples/0_Intro/module_api/Makefile @@ -11,8 +11,8 @@ ifeq (${HIP_PLATFORM}, hcc) all: runKernel.hip.out -vcpy_isa.compile: vcpy_isa.cpp - $(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_isa.cpp -o vcpy_isa.co +vcpy_kernel.compile: vcpy_kernel.cpp + $(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_kernel.cpp -o vcpy_kernel.co clean: rm -f *.co *.out @@ -21,15 +21,15 @@ endif ifeq (${HIP_PLATFORM}, nvcc) -vcpy_isa.compile: vcpy_isa.cu - $(HIPCC) --genco vcpy_isa.cu -o vcpy_isa.ptx +vcpy_kernel.compile: vcpy_kernel.cpp + $(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx clean: rm -f *.ptx *.out endif -all: vcpy_isa.compile runKernel.hip.out +all: vcpy_kernel.compile runKernel.hip.out runKernel.hip.out: runKernel.cpp $(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out diff --git a/projects/hip/samples/0_Intro/module_api/runKernel.cpp b/projects/hip/samples/0_Intro/module_api/runKernel.cpp index dcbaa4cd35..1156b8ddbf 100644 --- a/projects/hip/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/runKernel.cpp @@ -27,12 +27,12 @@ THE SOFTWARE. #define SIZE LEN<<2 #ifdef __HIP_PLATFORM_HCC__ -#define fileName "vcpy_isa.co" +#define fileName "vcpy_kernel.co" #define kernel_name "ZN12_GLOBAL__N_146_Z11hello_world16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf" #endif #ifdef __HIP_PLATFORM_NVCC__ -#define fileName "vcpy_isa.ptx" +#define fileName "vcpy_kernel.ptx" #define kernel_name "hello_world" #endif @@ -67,8 +67,8 @@ int main(){ hipModuleGetFunction(&Function, Module, kernel_name); #ifdef __HIP_PLATFORM_HCC__ - uint32_t len = LEN; - uint32_t one = 1; + uint32_t len = LEN; + uint32_t one = 1; std::vectorargBuffer(5); uint32_t *ptr32_t = (uint32_t*)&argBuffer[0]; 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.cpp b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp similarity index 91% rename from projects/hip/samples/0_Intro/module_api/vcpy_isa.cpp rename to projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp index ead3253115..640cf5b1bb 100644 --- a/projects/hip/samples/0_Intro/module_api/vcpy_isa.cpp +++ b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp @@ -6,4 +6,3 @@ __global__ void hello_world(hipLaunchParm lp, float *a, float *b) b[tx] = a[tx]; } -int main(){} From f31bc90537acdc6b1939629605e21c7e74c5b8d3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 2 Sep 2016 22:23:13 +0300 Subject: [PATCH 30/39] [HIPIFY] CUDA Driver API porting to HIP : populate Replacer's rules with enum members. enum members for: CUdevice_attribute/hipDeviceAttribute_t CUfunc_cache/hipFuncCache CUsharedconfig/hipSharedMemConfig [ROCm/hip commit: 8bc7aba7197b7f44860ce7957ace6f4689f897f2] --- projects/hip/clang-hipify/src/Cuda2Hip.cpp | 70 ++++++++++++++++------ 1 file changed, 52 insertions(+), 18 deletions(-) diff --git a/projects/hip/clang-hipify/src/Cuda2Hip.cpp b/projects/hip/clang-hipify/src/Cuda2Hip.cpp index 474f1a50dc..8959c1bc95 100644 --- a/projects/hip/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/hip/clang-hipify/src/Cuda2Hip.cpp @@ -247,35 +247,68 @@ struct cuda2hipMap { ///////////////////////////// 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"] = {"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["CUdevice_attribute_enum"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER}; + cuda2hipRename["CUdevice_attribute"] = {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER}; - cuda2hipRename["CUdevprop_st"] = {"hipDeviceProp_t", CONV_TYPE, API_DRIVER}; - cuda2hipRename["CUdevprop"] = {"hipDeviceProp_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["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}; + 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["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["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["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}; + 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 @@ -416,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}; From 85e7d1799b04be67a8836bca41cb73c0d0cc5eb1 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 2 Sep 2016 17:24:16 -0500 Subject: [PATCH 31/39] add hipStreamQuery on nvcc, fix hipStreamWaitEvent for nvcc [ROCm/hip commit: 1f40a73ae1f3719446f242baf9e4a4cd7cb780f7] --- projects/hip/include/nvcc_detail/hip_runtime_api.h | 5 +++++ .../hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp | 1 + 2 files changed, 6 insertions(+) diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index e2bcf41dd9..8e9b0d92a7 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -489,6 +489,11 @@ inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event 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) { diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 7bdae6a96f..c5a74b2bc0 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -21,6 +21,7 @@ THE SOFTWARE. #include "hip_runtime.h" #include "test_common.h" +#include unsigned p_streams = 6; From c3948922274d79e3f8f2c1081787c056d84a9d2a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 2 Sep 2016 17:50:31 -0500 Subject: [PATCH 32/39] remove printf Change-Id: I36510a7ff0425f05a6d0b9c65cda4397c9cc2d95 [ROCm/hip commit: 64bc513df98c39a4fa001b48b31bbb50f685a9c5] --- projects/hip/src/hip_stream.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index e044b6b578..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); @@ -95,7 +95,6 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int #if USE_AV_COPY if (stream != hipStreamNull) { - printf ("HIP: wait locked stream\n"); stream->locked_waitEvent(event); fastWait = true; // don't use the slow host-side synchronization. From e9f4f40a459710bbd6e484deb0b647587dbd1114 Mon Sep 17 00:00:00 2001 From: pensun Date: Sat, 3 Sep 2016 21:06:58 -0500 Subject: [PATCH 33/39] For module_api sample, use vcpy_kernel.cu to generate ptx file for NV path. Change-Id: Id0033678834288c4eaa56b12e7d447119be99deb [ROCm/hip commit: e08f4513c18c1546b3d8c982704e75a9dff2f782] --- projects/hip/samples/0_Intro/module_api/Makefile | 4 ++-- projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu | 6 ++++++ 2 files changed, 8 insertions(+), 2 deletions(-) create mode 100644 projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu diff --git a/projects/hip/samples/0_Intro/module_api/Makefile b/projects/hip/samples/0_Intro/module_api/Makefile index 99cff6bc0a..2570290f07 100644 --- a/projects/hip/samples/0_Intro/module_api/Makefile +++ b/projects/hip/samples/0_Intro/module_api/Makefile @@ -21,8 +21,8 @@ endif ifeq (${HIP_PLATFORM}, nvcc) -vcpy_kernel.compile: vcpy_kernel.cpp - $(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx +vcpy_kernel.compile: vcpy_kernel.cu + $(HIPCC) --genco vcpy_kernel.cu -o vcpy_kernel.ptx clean: rm -f *.ptx *.out diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu new file mode 100644 index 0000000000..ead7d70311 --- /dev/null +++ b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu @@ -0,0 +1,6 @@ +extern "C" __global__ void hello_world(float *a, float *b) +{ + int tx = threadIdx.x; + b[tx] = a[tx]; + +} From 576add359667b492f33bbae80a49fc1fbb0f4628 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 12:35:08 +0530 Subject: [PATCH 34/39] module_api: workaround to use vcpy_kernel.cpp for NV path Change-Id: Ib4868bf02c64070e846c19427c39289609909466 [ROCm/hip commit: 07badd6f4e9567f2503abf61431cc0771752f464] --- .../hip/samples/0_Intro/module_api/Makefile | 19 ++++++++----------- .../samples/0_Intro/module_api/runKernel.cpp | 9 ++++++--- .../0_Intro/module_api/vcpy_kernel.cpp | 8 +++++++- .../samples/0_Intro/module_api/vcpy_kernel.cu | 6 ------ 4 files changed, 21 insertions(+), 21 deletions(-) delete mode 100644 projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu diff --git a/projects/hip/samples/0_Intro/module_api/Makefile b/projects/hip/samples/0_Intro/module_api/Makefile index 2570290f07..bda0a7dd93 100644 --- a/projects/hip/samples/0_Intro/module_api/Makefile +++ b/projects/hip/samples/0_Intro/module_api/Makefile @@ -6,13 +6,15 @@ HIPCC=$(HIP_PATH)/bin/hipcc OPT= HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) -ifeq (${HIP_PLATFORM}, hcc) - GENCODEOBJECT_FLAGS=--target-isa=fiji +all: vcpy_kernel.compile runKernel.hip.out -all: runKernel.hip.out +runKernel.hip.out: runKernel.cpp + $(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out + +ifeq (${HIP_PLATFORM}, hcc) vcpy_kernel.compile: vcpy_kernel.cpp - $(HIPCC) --genco $(GENCODEOBJECT_FLAGS) vcpy_kernel.cpp -o vcpy_kernel.co + $(HIPCC) --genco --target-isa=fiji vcpy_kernel.cpp -o vcpy_kernel.co clean: rm -f *.co *.out @@ -21,16 +23,11 @@ endif ifeq (${HIP_PLATFORM}, nvcc) -vcpy_kernel.compile: vcpy_kernel.cu - $(HIPCC) --genco vcpy_kernel.cu -o vcpy_kernel.ptx +vcpy_kernel.compile: vcpy_kernel.cpp + $(HIPCC) --genco vcpy_kernel.cpp -o vcpy_kernel.ptx clean: rm -f *.ptx *.out endif -all: vcpy_kernel.compile runKernel.hip.out - -runKernel.hip.out: runKernel.cpp - $(HIPCC) $(OPT) runKernel.cpp -o runKernel.hip.out - diff --git a/projects/hip/samples/0_Intro/module_api/runKernel.cpp b/projects/hip/samples/0_Intro/module_api/runKernel.cpp index 1156b8ddbf..ea372773c3 100644 --- a/projects/hip/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/runKernel.cpp @@ -83,9 +83,12 @@ int main(){ #endif #ifdef __HIP_PLATFORM_NVCC__ - std::vectorargBuffer(2); - memcpy(&argBuffer[0], &Ad, sizeof(void*)); - memcpy(&argBuffer[1], &Bd, sizeof(void*)); + 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 diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp index 640cf5b1bb..ebb6d066a9 100644 --- a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp @@ -1,6 +1,12 @@ #include -__global__ void hello_world(hipLaunchParm lp, float *a, float *b) +#ifdef __HIP_PLATFORM_NVCC__ +#define EXTERN_C extern "C" +#else +#define EXTERN_C +#endif + +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/samples/0_Intro/module_api/vcpy_kernel.cu b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cu deleted file mode 100644 index ead7d70311..0000000000 --- a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.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]; - -} From 9f2206c8b1868aa3721817e90f71ab4639348257 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 13:48:44 +0530 Subject: [PATCH 35/39] hipgenisa.sh: handle extern C on HCC path Change-Id: I2daf09d2cb323ab34653988b668e8bd89f6db94a [ROCm/hip commit: fcd401dfb759629afea50f4231aeb688cc23e8ff] --- projects/hip/bin/hipgenisa.sh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip/bin/hipgenisa.sh b/projects/hip/bin/hipgenisa.sh index 03952f49d8..aba3d48e8c 100755 --- a/projects/hip/bin/hipgenisa.sh +++ b/projects/hip/bin/hipgenisa.sh @@ -22,7 +22,8 @@ HIP_PATH="$( command cd -P "$( dirname "$SOURCE" )/.." && pwd )" export KMDUMPISA=1 export KMDUMPLLVM=1 mkdir /tmp/hipgenisa -cp $FILE_NAMES $FILE_NAMES.kernel.tmp.cpp +#cp $FILE_NAMES $FILE_NAMES.kernel.tmp.cpp +sed 's/extern \+"C" \+//g' $FILE_NAMES > $FILE_NAMES.kernel.tmp.cpp echo " int main(){} " >> $FILE_NAMES.kernel.tmp.cpp From 6155170faf5e854cb93ee59c3cb656db6ecd54e8 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 13:49:43 +0530 Subject: [PATCH 36/39] module_api sample: no longer need EXTERN_C workaround Change-Id: Ida087d832df8e1f3620b38f920ec2853aad641c8 [ROCm/hip commit: ba8a05a3c0e072a37b434f32a263c78c872661ed] --- .../hip/samples/0_Intro/module_api/vcpy_kernel.cpp | 10 ++-------- 1 file changed, 2 insertions(+), 8 deletions(-) diff --git a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp index ebb6d066a9..0e051f76fc 100644 --- a/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/vcpy_kernel.cpp @@ -1,12 +1,6 @@ -#include +#include "hip/hip_runtime.h" -#ifdef __HIP_PLATFORM_NVCC__ -#define EXTERN_C extern "C" -#else -#define EXTERN_C -#endif - -EXTERN_C __global__ void hello_world(hipLaunchParm lp, float *a, float *b) +extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b) { int tx = hipThreadIdx_x; b[tx] = a[tx]; From a9973913e083f40b37368a312f2e6be095cca771 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 16:20:49 +0530 Subject: [PATCH 37/39] hipgenisa.sh improvements - hipgenisa now unmangles kernel names - hipgenisa now uses a random temporary directory Change-Id: I1d0204336bd8042dcf3f39c8f905fcceebd6b785 [ROCm/hip commit: fbc726d1b97e87915597b96f5282ac26b437af27] --- projects/hip/bin/hipgenisa.sh | 23 ++++++++++++++--------- 1 file changed, 14 insertions(+), 9 deletions(-) diff --git a/projects/hip/bin/hipgenisa.sh b/projects/hip/bin/hipgenisa.sh index aba3d48e8c..f52b1c5f4e 100755 --- a/projects/hip/bin/hipgenisa.sh +++ b/projects/hip/bin/hipgenisa.sh @@ -21,17 +21,22 @@ HIP_PATH="$( command cd -P "$( dirname "$SOURCE" )/.." && pwd )" export KMDUMPISA=1 export KMDUMPLLVM=1 -mkdir /tmp/hipgenisa -#cp $FILE_NAMES $FILE_NAMES.kernel.tmp.cpp +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 /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 +$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 /tmp/hipgenisa -export KMDUMPISA=0 -export KMDUMPLLVM=0 +rm -r $hipgenisa_dir From 301ae0fa90bc9e38b4d77b25f656f702ee8457cd Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 16:25:02 +0530 Subject: [PATCH 38/39] hip_kernel_language.md: Update documentation for kernel compilation Change-Id: I51073363c3e8a8837c5fe2ae5a04dae0d441229e [ROCm/hip commit: b8495e5961ef6826b6999c48696eff95878c704a] --- projects/hip/docs/markdown/hip_kernel_language.md | 9 --------- 1 file changed, 9 deletions(-) diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index c9df5ce1b1..628cd36c26 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -678,14 +678,5 @@ The file format for binary is `.co` which means Code Object. The following comma [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. - -To load a kernel into HIP, we need both the code object and the name of the kernel stored within the code object. -In order to get name of the kernel, use: -``` -$ objdump -x [CODE_OBJECT_FILE]`. -``` -CODE_OBJECT_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++. The hipLoadKernelModule API needs to specify this mangled name on the HIP/hcc path. - From 7776ff3ecbb6a2c865e0749696f20a345f446339 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 4 Sep 2016 16:26:16 +0530 Subject: [PATCH 39/39] module_api: HCC path no longer needs mangled kernel name Change-Id: I4c1cb218bfdd05c9fba57276167e3e4205b93614 [ROCm/hip commit: 7020584cdf806aa846426bd8dc4e902e2b806f96] --- projects/hip/samples/0_Intro/module_api/runKernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/samples/0_Intro/module_api/runKernel.cpp b/projects/hip/samples/0_Intro/module_api/runKernel.cpp index ea372773c3..1c0dd9f349 100644 --- a/projects/hip/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip/samples/0_Intro/module_api/runKernel.cpp @@ -28,7 +28,7 @@ THE SOFTWARE. #ifdef __HIP_PLATFORM_HCC__ #define fileName "vcpy_kernel.co" -#define kernel_name "ZN12_GLOBAL__N_146_Z11hello_world16grid_launch_parmPfS0__functor19__cxxamp_trampolineEiiiiiiPKfPf" +#define kernel_name "hello_world" #endif #ifdef __HIP_PLATFORM_NVCC__