From cdf0813551ac561821dd7232816562762ccbd860 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 3 Apr 2017 20:32:50 +0300 Subject: [PATCH 01/19] [HIPIFY] Remove hipLaunchParm in HIP kernel declaration. [ROCm/hip commit: 449d21b3c6c3cca098e51fa0389de55ccb26b015] --- projects/hip/hipify-clang/src/Cuda2Hip.cpp | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp index 6c24fbf288..27447b8d8c 100644 --- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp @@ -2343,9 +2343,6 @@ private: LangOptions DefaultLangOptions; SmallString<40> XStr; raw_svector_ostream OS(XStr); - StringRef initialParamList; - OS << "hipLaunchParm lp"; - size_t repLength = OS.str().size(); SourceLocation sl = kernelDecl->getNameInfo().getEndLoc(); SourceLocation kernelArgListStart = Lexer::findLocationAfterToken(sl, tok::l_paren, *SM, DefaultLangOptions, true); DEBUG(dbgs() << kernelArgListStart.printToString(*SM)); @@ -2355,14 +2352,12 @@ private: SourceLocation kernelArgListStart(pvdFirst->getLocStart()); SourceLocation kernelArgListEnd(pvdLast->getLocEnd()); SourceLocation stop = Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); - repLength += SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); - initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), repLength); - OS << ", " << initialParamList; + size_t repLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); + OS << StringRef(SM->getCharacterData(kernelArgListStart), repLength); + Replacement Rep0(*(Result.SourceManager), kernelArgListStart, repLength, OS.str()); + FullSourceLoc fullSL(sl, *(Result.SourceManager)); + insertReplacement(Rep0, fullSL); } - DEBUG(dbgs() << "initial paramlist: " << initialParamList << "\n" << "new paramlist: " << OS.str() << "\n"); - Replacement Rep0(*(Result.SourceManager), kernelArgListStart, repLength, OS.str()); - FullSourceLoc fullSL(sl, *(Result.SourceManager)); - insertReplacement(Rep0, fullSL); } bool cudaCall(const MatchFinder::MatchResult &Result) { From b203f9f5939ed79f565d98077ba378a16aedf08f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 3 Apr 2017 22:05:01 +0300 Subject: [PATCH 02/19] [HIPIFY] GGL support + hipLaunchKernel -> hipLaunchKernelGGL + macro HIP_KERNEL_NAME is no longer used [ROCm/hip commit: 143e0af716b279f3a1e15b0ef1bf62103ddcd0a7] --- projects/hip/hipify-clang/src/Cuda2Hip.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp index 27447b8d8c..383af0440c 100644 --- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp @@ -2426,9 +2426,9 @@ private: XStr.clear(); if (calleeName.find(',') != StringRef::npos) { SmallString<128> tmpData; - calleeName = Twine("HIP_KERNEL_NAME(" + calleeName + ")").toStringRef(tmpData); + calleeName = Twine("(" + calleeName + ")").toStringRef(tmpData); } - OS << "hipLaunchKernel(" << calleeName << ","; + OS << "hipLaunchKernelGGL(" << calleeName << ","; const CallExpr *config = launchKernel->getConfig(); DEBUG(dbgs() << "Kernel config arguments:" << "\n"); SourceManager *SM = Result.SourceManager; @@ -2468,7 +2468,7 @@ private: Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str()); FullSourceLoc fullSL(launchKernel->getLocStart(), *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernel", CONV_KERN, API_RUNTIME}; + hipCounter counter = {"hipLaunchKernelGGL", CONV_KERN, API_RUNTIME}; updateCounters(counter, refName.str()); return true; } From 564a90cb317a4a6cd0cd93ea72a5d1dacecc18a1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 4 Apr 2017 08:06:09 +0530 Subject: [PATCH 03/19] hip_doc packaging script updates - Gracefully handle missing doxygen and grip tools Change-Id: I1a4a653d687c136c6d9237062ab4d02bc6cb3db1 [ROCm/hip commit: 958af4dae14caa91f5f3a5efee5de3d78e3c9b78] --- projects/hip/packaging/hip_doc.txt | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/projects/hip/packaging/hip_doc.txt b/projects/hip/packaging/hip_doc.txt index bbcaf54ec8..d5a0c471b1 100644 --- a/projects/hip/packaging/hip_doc.txt +++ b/projects/hip/packaging/hip_doc.txt @@ -1,12 +1,19 @@ cmake_minimum_required(VERSION 2.8.3) project(hip_doc) -add_custom_target(build_doxygen ALL +find_program(DOXYGEN_EXE doxygen) +if(DOXYGEN_EXE) + add_custom_target(build_doxygen ALL COMMAND HIP_PATH=@hip_SOURCE_DIR@ doxygen @hip_SOURCE_DIR@/docs/doxygen-input/doxy.cfg) -add_custom_target(convert_md_to_html ALL + install(DIRECTORY RuntimeAPI/html DESTINATION docs/docs/RuntimeAPI) +endif() + +find_program(GRIP_EXE grip) +if(GRIP_EXE) + add_custom_target(convert_md_to_html ALL COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html) -install(DIRECTORY RuntimeAPI/html DESTINATION docs/docs/RuntimeAPI) -install(DIRECTORY md2html/ DESTINATION docs) + install(DIRECTORY md2html/ DESTINATION docs) +endif() ############################# # Packaging steps From 013a98ad2bc5d5ef839117dfc341b452e195a97b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 4 Apr 2017 08:07:56 +0530 Subject: [PATCH 04/19] Remove unused packaging scripts Change-Id: I609391b905810eb24f7fd4ea6d7f27166ca001b3 [ROCm/hip commit: 93887e9c83d883d9cb8c92cce3dc7356b9f25a56] --- .../packaging/create_hip_samples_installer.sh | 23 ------------------- 1 file changed, 23 deletions(-) delete mode 100755 projects/hip/packaging/create_hip_samples_installer.sh diff --git a/projects/hip/packaging/create_hip_samples_installer.sh b/projects/hip/packaging/create_hip_samples_installer.sh deleted file mode 100755 index 91789d2524..0000000000 --- a/projects/hip/packaging/create_hip_samples_installer.sh +++ /dev/null @@ -1,23 +0,0 @@ -#!/bin/bash -function die { - echo "${1-Died}." >&2 - exit 1 -} - -payload=$1 -script=$2 -[ "$payload" != "" ] || [ "$script" != "" ] || die "Invalid arguments!" -tmp=__extract__$RANDOM - -printf "#!/bin/bash -samples_dir=\$1 -[ \"\$samples_dir\" != \"\" ] || read -e -p \"Enter the path to extract the HIP samples: \" samples_dir -mkdir -p \$samples_dir -PAYLOAD=\`awk '/^__PAYLOAD_BELOW__/ {print NR + 1; exit 0; }' \$0\` -tail -n+\$PAYLOAD \$0 | tar -xz -C \$samples_dir -echo \"HIP samples installed in \$samples_dir\" -exit 0 -__PAYLOAD_BELOW__\n" > "$tmp" - -cat "$tmp" "$payload" > "$script" && rm "$tmp" -chmod +x "$script" From 06b32795c81e798454ea8ddd0cdb09796dbd1dfa Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 4 Apr 2017 15:51:10 +0530 Subject: [PATCH 05/19] mgpu IPC support fix Change-Id: I12e4b2fd189c3658efd3b07defa18ece3853b0eb [ROCm/hip commit: fc61b793fe0e53896509f71b8c3b1c401aa7312d] --- projects/hip/src/hip_memory.cpp | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 805fc9efc0..da5530349f 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1260,10 +1260,15 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*) &handle; //Attach ipc memory - hsa_status_t hsa_status = - hsa_amd_ipc_memory_attach((hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, 1, agent, devPtr); - if(hsa_status != HSA_STATUS_SUCCESS) - hipStatus = hipErrorMapBufferObjectFailed; + auto ctx= ihipGetTlsDefaultCtx(); + { + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); + // the peerCnt always stores self so make sure the trace actually + hsa_status_t hsa_status = + hsa_amd_ipc_memory_attach((hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(), crit->peerAgents(), devPtr); + if(hsa_status != HSA_STATUS_SUCCESS) + hipStatus = hipErrorMapBufferObjectFailed; + } #else hipStatus = hipErrorRuntimeOther; #endif From 735c6df1bc03030f1247f434ae0d4c230c8a53a1 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Wed, 5 Apr 2017 16:23:27 -0500 Subject: [PATCH 06/19] hip_debug.md update from Alex, regarding __device__ function restrictions Change-Id: I5e54fd97fc632d4283f76282e3935396a1aad235 [ROCm/hip commit: a98802f72c40f96a3ad23cc4902417ca17799f2c] --- projects/hip/docs/markdown/hip_bugs.md | 90 +++++++++++++++++++++++++- 1 file changed, 89 insertions(+), 1 deletion(-) diff --git a/projects/hip/docs/markdown/hip_bugs.md b/projects/hip/docs/markdown/hip_bugs.md index e15c37fc54..addf2c17f6 100644 --- a/projects/hip/docs/markdown/hip_bugs.md +++ b/projects/hip/docs/markdown/hip_bugs.md @@ -5,6 +5,7 @@ - [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#errors-related-to-undefined-reference-to-hclaunchkernel__grid_launch_parm) - [Application hangs after a hipLaunchKernel call](#what-if-i-see-application-hangs-after-a-hiplaunchkernel-call) - [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method) +- [HIP is more restrictive in enforcing restrictions](#hip-is-more-restrictive-in-enforcing-restrictions) @@ -46,4 +47,91 @@ hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_H ### What is the current limitation of HIP Generic Grid Launch method? 1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak))); -2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch. \ No newline at end of file +2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch. + + +### HIP is more restrictive in enforcing restrictions +By the language specification, both for HIP and CUDA it is forbidden to call a +`__device__` function in a `__host__` context. In practice, you may observe +differences in the strictness of this restriction, with HIP exhibiting a tighter +adherence to the specification i.e. being less tolerant of infringing code. The +solution is to always ensure that all functions which are to be called in a +`__device__` context are correctly annotated to reflect it. An interesting case +where these differences emerge is shown below (this has been lifted from +production code, and relies on a the common [C++ Member Detector idiom][1], as it +would be implemented pre C++11): +```c++ +#include +#include + +struct meta_yes { char a[1]; }; +struct meta_no { char a[2]; }; + +// Dual restriction is necessary in HIP if the detector is to work for +// __device__ contexts as well as __host__ ones. NVCC is less strict. +template +__host__ __device__ +const T& return_ref(); + +template +struct has_nullary_operator { + // Dual restriction is necessary in HIP if the detector is to work for + // __device__ contexts as well as __host__ ones. NVCC is less strict. + template + __host__ __device__ + static + meta_yes testFunctor( + C const *, + typename std::enable_if< + (sizeof(return_ref().operator()()) > 0)>::type* = nullptr); + static + meta_no testFunctor(...); + + enum { + value = sizeof(testFunctor(static_cast(0))) == sizeof(meta_yes) }; +}; + +template< + typename Scalar, + typename NullaryOp, + bool has_nullary = has_nullary_operator::value> +struct nullary_wrapper { + template + T packetOp() const { return T{1}; } +}; + + +template +struct nullary_wrapper { + template + T packetOp() const { return T{10}; } +}; + +// This specialisation will fail to compile. +template +struct nullary_wrapper {}; + +template +struct UniformRandomGenerator; + +template<> struct UniformRandomGenerator { + float operator()() const [[hc]] { return 42.0; } +}; + +__device__ +void this_will_not_compile_if_detector_is_not_marked_device() +{ + float f = + nullary_wrapper< + float, UniformRandomGenerator>().packetOp(); +} + +__host__ +void this_will_not_compile_if_detector_is_marked_device_only() +{ + float f = + nullary_wrapper< + float, UniformRandomGenerator>().packetOp(); +} +``` +[1]: https://en.wikibooks.org/wiki/More_C%2B%2B_Idioms/Member_Detector \ No newline at end of file From d957b0f562c12d342e5ea13133419c159e7e2f42 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 5 Apr 2017 17:43:20 -0500 Subject: [PATCH 07/19] Add bug descrip for "no matching constructor" [ROCm/hip commit: b728637692461cdc08e7af04d06644ab79d19f94] --- projects/hip/docs/markdown/hip_bugs.md | 66 +++++++++++++++++++++++++- 1 file changed, 65 insertions(+), 1 deletion(-) diff --git a/projects/hip/docs/markdown/hip_bugs.md b/projects/hip/docs/markdown/hip_bugs.md index addf2c17f6..234dec4e0e 100644 --- a/projects/hip/docs/markdown/hip_bugs.md +++ b/projects/hip/docs/markdown/hip_bugs.md @@ -50,6 +50,70 @@ hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_H 2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch. +### Errors related to `no matching constructor` + +The symptom is the compiler would complain about errors like `no matching constructor` for classes/structs passed as arguments into a GPU kernel. Often, this is caused by a design limitation in HCC where array-typed member variables inside a class/struct can’t be correctly passed into GPU kernels. To mitigate this issue, a custom serializer/deserializer pair is provided. + +For example, `Foo` in the code snippets below contains an array-typed member variable `table`, which would fail the compiler if used as a kernel argument. + +``` +struct Foo { + // table is an array, which makes foo + int table[3]; +}; +``` + +An workaround is to provide a custom serializer on CPU side, and append the contents of the array as kernel arguments: + +``` + +struct Foo { + int table[3]; + + // user-provided CPU serializer + // must append the contents of the array member as kernel arguments +#ifdef __HCC__ + __attribute__((annotate(“serialize”))) + void __cxxamp_serialize(Kalmar::Serialize &s) const { + for (int i = 0; i < 3; ++i) + s.Append(sizeof(int), &table[i]); + } +#endif +}; +``` + +Then, provide a custom deserializer on GPU side, to help reconstruct the array within GPU kernels. Notice that the deserializer can not be a function template, and should have scalar-typed parameters of the number equals to the length of the array-typed member variable. For example: + +``` +struct Foo { + int table[3]; + + // user-provided GPU deserializer + // table has 3 int elements, so deserializer must have 3 int parameters. +#ifdef __HCC__ + __attribute__((annotate(“user_deserialize”))) + Foo(int x0, int x1, int x2) [[cpu]][[hc]] { + table[0] = x0; + table[1] = x1; + table[2] = x2; + } +#endif + +#ifdef __HCC__ + __attribute__((annotate(“serialize”))) + void __cxxamp_serialize(Kalmar::Serialize &s) const { + s.Append(sizeof(int), &table[0]); + s.Append(sizeof(int), &table[1]); + s.Append(sizeof(int), &table[2]); + } +#endif +}; +``` + + +Rather than create serializer functions, another workaround is to pass the member fields from the structure as simple data types. + + ### HIP is more restrictive in enforcing restrictions By the language specification, both for HIP and CUDA it is forbidden to call a `__device__` function in a `__host__` context. In practice, you may observe @@ -134,4 +198,4 @@ void this_will_not_compile_if_detector_is_marked_device_only() float, UniformRandomGenerator>().packetOp(); } ``` -[1]: https://en.wikibooks.org/wiki/More_C%2B%2B_Idioms/Member_Detector \ No newline at end of file +[1]: https://en.wikibooks.org/wiki/More_C%2B%2B_Idioms/Member_Detector From dab8a2d0876a9bb7cb5616671c5f420f457d6c71 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 5 Apr 2017 19:40:00 -0500 Subject: [PATCH 08/19] Doc cleanup & add bug descript for restrict specifier issue [ROCm/hip commit: aaadc6026457d207db223b151ef1b2f2dceca737] --- projects/hip/docs/markdown/hip_bugs.md | 92 ++++++++++---------------- 1 file changed, 34 insertions(+), 58 deletions(-) diff --git a/projects/hip/docs/markdown/hip_bugs.md b/projects/hip/docs/markdown/hip_bugs.md index 234dec4e0e..c53b68d796 100644 --- a/projects/hip/docs/markdown/hip_bugs.md +++ b/projects/hip/docs/markdown/hip_bugs.md @@ -15,9 +15,11 @@ Some common code practices may lead to hipcc generating a error with the form : undefined reference to `__hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T_ To workaround, try: -- Avoid calling hcLaunchKernel from a function with the __host__ attribute +- Avoid calling hipLaunchKernel from a function with the __host__ attribute +``` __host__ MyFunc(…) { hipLaunchKernel(myKernel, …) +``` - Avoid use of static with kernel definition: static __global__ MyKernel - Avoid defining kernels in anonymous namespace @@ -25,25 +27,6 @@ namespace { __global__ MyKernel … - Avoid calling member functions -If hipLaunchKernel takes parameters that request explicitly memcpy, then it will cause application hang. -Reason is that the hipLaunchKernel macro locks the stream. -If kernel paramters are actually function calls which invoke other hip apis (i.e. memcpy) to the same stream, then deadlock occurs. - -To workaround, try: -Move the function calls so they occur outside the hipLaunchKernel macro, store results in temps, then use the tems inside the kernel. - -``` -// Example pseudo code causing system hang: -// "bottom[0]->gpu_data()" calls hipMemcpy() implicitly and using the same stream, cause deadlock condition. -hipLaunchKernel(HIP_KERNEL_NAME(LRNComputeDiff),dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads, - bottom[0]->gpu_data()); - -// Move "gpu_data()" ouside of hipLaunchKernel to avoid hang. -auto bot_gpu_data = bottom[0]->gpu_data(); -hipLaunchKernel( LRNComputeDiff, dim3(CAFFE_GET_BLOCKS(n_threads)), dim3(CAFFE_HIP_NUM_THREADS), 0, 0, n_threads, - bot_gpu_data); - -``` ### What is the current limitation of HIP Generic Grid Launch method? 1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak))); @@ -115,87 +98,80 @@ Rather than create serializer functions, another workaround is to pass the membe ### HIP is more restrictive in enforcing restrictions -By the language specification, both for HIP and CUDA it is forbidden to call a +The language specification for HIP and CUDA forbid calling a `__device__` function in a `__host__` context. In practice, you may observe differences in the strictness of this restriction, with HIP exhibiting a tighter -adherence to the specification i.e. being less tolerant of infringing code. The -solution is to always ensure that all functions which are to be called in a +adherence to the specification and thus less tolerant of infringing code. The +solution is to ensure that all functions which are called in a `__device__` context are correctly annotated to reflect it. An interesting case -where these differences emerge is shown below (this has been lifted from -production code, and relies on a the common [C++ Member Detector idiom][1], as it -would be implemented pre C++11): +where these differences emerge is shown below. This relies on a the common +[C++ Member Detector idiom][1], as it would be implemented pre C++11): + ```c++ #include #include -struct meta_yes { char a[1]; }; -struct meta_no { char a[2]; }; +struct aye { bool a[1]; }; +struct nay { bool a[2]; }; // Dual restriction is necessary in HIP if the detector is to work for // __device__ contexts as well as __host__ ones. NVCC is less strict. template __host__ __device__ -const T& return_ref(); +const T& cref_t(); template -struct has_nullary_operator { +struct Has_call_operator { // Dual restriction is necessary in HIP if the detector is to work for // __device__ contexts as well as __host__ ones. NVCC is less strict. template __host__ __device__ static - meta_yes testFunctor( + aye test( C const *, typename std::enable_if< - (sizeof(return_ref().operator()()) > 0)>::type* = nullptr); + (sizeof(cref_t().operator()()) > 0)>::type* = nullptr); static - meta_no testFunctor(...); + nay test(...); - enum { - value = sizeof(testFunctor(static_cast(0))) == sizeof(meta_yes) }; + enum { value = sizeof(test(static_cast(0))) == sizeof(aye) }; }; -template< - typename Scalar, - typename NullaryOp, - bool has_nullary = has_nullary_operator::value> -struct nullary_wrapper { - template - T packetOp() const { return T{1}; } +template::value> +struct Wrapper { + template + V f() const { return T{1}; } }; -template -struct nullary_wrapper { - template - T packetOp() const { return T{10}; } +template +struct Wrapper { + template + V f() const { return T{10}; } }; -// This specialisation will fail to compile. -template -struct nullary_wrapper {}; +// This specialisation will yield a compile-time error, if selected. +template +struct Wrapper {}; template -struct UniformRandomGenerator; +struct Functor; -template<> struct UniformRandomGenerator { - float operator()() const [[hc]] { return 42.0; } +template<> struct Functor { + __device__ + float operator()() const { return 42.0f; } }; __device__ void this_will_not_compile_if_detector_is_not_marked_device() { - float f = - nullary_wrapper< - float, UniformRandomGenerator>().packetOp(); + float f = Wrapper>().f(); } __host__ void this_will_not_compile_if_detector_is_marked_device_only() { - float f = - nullary_wrapper< - float, UniformRandomGenerator>().packetOp(); + float f = Wrapper>().f(); } ``` [1]: https://en.wikibooks.org/wiki/More_C%2B%2B_Idioms/Member_Detector From 8d306c0578e725949ea43dc3db70ab04b64f72ee Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 5 Apr 2017 21:59:11 -0500 Subject: [PATCH 09/19] Update bug workarounds to reflect tool improvements. [ROCm/hip commit: 00bf446d53b7def0433d52a99c0cef646655e3d9] --- projects/hip/docs/markdown/hip_bugs.md | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/projects/hip/docs/markdown/hip_bugs.md b/projects/hip/docs/markdown/hip_bugs.md index c53b68d796..14f2935f17 100644 --- a/projects/hip/docs/markdown/hip_bugs.md +++ b/projects/hip/docs/markdown/hip_bugs.md @@ -9,23 +9,23 @@ -### Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm** +### Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**` Some common code practices may lead to hipcc generating a error with the form : undefined reference to `__hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T_ -To workaround, try: -- Avoid calling hipLaunchKernel from a function with the __host__ attribute -``` -__host__ MyFunc(…) { -hipLaunchKernel(myKernel, …) -``` +Suggested workarounds: - Avoid use of static with kernel definition: +```c++ static __global__ MyKernel -- Avoid defining kernels in anonymous namespace +``` + +- Avoid defining kernels in anonymous namespace : +```c++ namespace { -__global__ MyKernel … -- Avoid calling member functions + __global__ MyKernel +} +``` ### What is the current limitation of HIP Generic Grid Launch method? From 72238e21191501981a4d0012ffcc823baed4f153 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Wed, 5 Apr 2017 22:25:41 -0500 Subject: [PATCH 10/19] add extra guard to grid_launch_GGL header Change-Id: I120619c08ea2d084804fcb1639efbe6c4648dde9 [ROCm/hip commit: 1a8ea824972c99c0c78fef15e91170dc85ecf949] --- projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp b/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp index 4fd7c3ff3a..1d765dfc48 100644 --- a/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp +++ b/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp @@ -21,6 +21,7 @@ THE SOFTWARE. */ #pragma once +#if GENERIC_GRID_LAUNCH == 1 #include "concepts.hpp" #include "helpers.hpp" @@ -851,3 +852,4 @@ namespace hip_impl ##__VA_ARGS__);\ } } +#endif //GENERIC_GRID_LAUNCH From 26f6ce992d99a4d9c098d6a4972dbfb46a7af0e1 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 6 Apr 2017 09:29:44 -0500 Subject: [PATCH 11/19] added __host__ for complex functions and corrected memset and memcpy test Change-Id: I9ffefb7a0025aa111a54d20d2766982df15532e7 [ROCm/hip commit: 6dff2714e95ac65c6662d390dbc57dd232627b66] --- .../hip/include/hip/hcc_detail/hip_complex.h | 50 +++++++++---------- .../tests/src/deviceLib/hipDeviceMemcpy.cpp | 45 ++++++++++++----- 2 files changed, 58 insertions(+), 37 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 9ff75d381a..dd742e484c 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -177,45 +177,45 @@ COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long) #endif -__device__ static inline float hipCrealf(hipFloatComplex z){ +__device__ __host__ static inline float hipCrealf(hipFloatComplex z){ return z.x; } -__device__ static inline float hipCimagf(hipFloatComplex z){ +__device__ __host__ static inline float hipCimagf(hipFloatComplex z){ return z.y; } -__device__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){ +__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){ hipFloatComplex z; z.x = a; z.y = b; return z; } -__device__ static inline hipFloatComplex hipConjf(hipFloatComplex z){ +__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z){ hipFloatComplex ret; ret.x = z.x; ret.y = -z.y; return ret; } -__device__ static inline float hipCsqabsf(hipFloatComplex z){ +__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z){ return z.x * z.x + z.y * z.y; } -__device__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){ +__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x + q.x, p.y + q.y); } -__device__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){ +__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x - q.x, p.y - q.y); } -__device__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){ +__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){ return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); } -__device__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ +__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){ float sqabs = hipCsqabsf(q); hipFloatComplex ret; ret.x = (p.x * q.x + p.y * q.y)/sqabs; @@ -223,51 +223,51 @@ __device__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatCom return ret; } -__device__ static inline float hipCabsf(hipFloatComplex z){ +__device__ __host__ static inline float hipCabsf(hipFloatComplex z){ return sqrtf(hipCsqabsf(z)); } -__device__ static inline double hipCreal(hipDoubleComplex z){ +__device__ __host__ static inline double hipCreal(hipDoubleComplex z){ return z.x; } -__device__ static inline double hipCimag(hipDoubleComplex z){ +__device__ __host__ static inline double hipCimag(hipDoubleComplex z){ return z.y; } -__device__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){ +__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){ hipDoubleComplex z; z.x = a; z.y = b; return z; } -__device__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){ +__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){ hipDoubleComplex ret; ret.x = z.x; ret.y = z.y; return ret; } -__device__ static inline double hipCsqabs(hipDoubleComplex z){ +__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z){ return z.x * z.x + z.y * z.y; } -__device__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){ +__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x + q.x, p.y + q.y); } -__device__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){ +__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x - q.x, p.y - q.y); } -__device__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){ +__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){ return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y); } -__device__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ +__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){ double sqabs = hipCsqabs(q); hipDoubleComplex ret; ret.x = (p.x * q.x + p.y * q.y)/sqabs; @@ -275,28 +275,28 @@ __device__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleC return ret; } -__device__ static inline double hipCabs(hipDoubleComplex z){ +__device__ __host__ static inline double hipCabs(hipDoubleComplex z){ return sqrtf(hipCsqabs(z)); } typedef hipFloatComplex hipComplex; -__device__ static inline hipComplex make_hipComplex(float x, +__device__ __host__ static inline hipComplex make_hipComplex(float x, float y){ return make_hipFloatComplex(x, y); } -__device__ static inline hipFloatComplex hipComplexDoubleToFloat +__device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat (hipDoubleComplex z){ return make_hipFloatComplex((float)z.x, (float)z.y); } -__device__ static inline hipDoubleComplex hipComplexFloatToDouble +__device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble (hipFloatComplex z){ return make_hipDoubleComplex((double)z.x, (double)z.y); } -__device__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ +__device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){ float real = (p.x * q.x) + r.x; float imag = (q.x * p.y) + r.y; @@ -306,7 +306,7 @@ __device__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComp return make_hipComplex(real, imag); } -__device__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){ +__device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){ float real = (p.x * q.x) + r.x; float imag = (q.x * p.y) + r.y; diff --git a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp index 54fd02c0c2..3843c07bb9 100644 --- a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp +++ b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp @@ -1,18 +1,29 @@ -#include +#include #include "hip/hip_runtime.h" #include "hip/hip_runtime_api.h" +#include "../test_common.h" + #define LEN 1030 #define SIZE LEN << 2 -__global__ void cpy(hipLaunchParm lp, uint32_t *Out, uint32_t *In, uint32_t *Vald) +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ + + +__global__ void cpy(hipLaunchParm lp, uint32_t *Out, uint32_t *In) { - memcpy(Out, In, SIZE, Vald); + int tx = hipThreadIdx_x; + memcpy(Out + tx, In + tx, SIZE/LEN); } __global__ void set(hipLaunchParm lp, uint32_t *ptr, uint8_t val, size_t size) { - memset(ptr, val, size); + int tx = hipThreadIdx_x; + memset(ptr + tx, val, size); } int main() @@ -24,19 +35,29 @@ int main() Val = new uint32_t; *Val = 0; for(int i=0;i Date: Thu, 6 Apr 2017 10:48:11 -0500 Subject: [PATCH 12/19] GGL update, add while 0 guard for hipLaunchKernel API Change-Id: Ie48ef8ca2ab5e26a51febfcd92417902c33fbf66 [ROCm/hip commit: ad18bf3b29a2bbad52bb54559949c04a84d18174] --- .../hip/hcc_detail/grid_launch_GGL.hpp | 21 ++++++++++--------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp b/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp index 1d765dfc48..8f1abbb70b 100644 --- a/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp +++ b/projects/hip/include/hip/hcc_detail/grid_launch_GGL.hpp @@ -841,15 +841,16 @@ namespace hip_impl group_mem_bytes,\ stream,\ ...)\ - {\ - hipLaunchKernelGGL(\ - kernel_name,\ - num_blocks,\ - dim_blocks,\ - group_mem_bytes,\ - stream,\ - hipLaunchParm{},\ - ##__VA_ARGS__);\ - } + do {\ + hipLaunchKernelGGL(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + hipLaunchParm{},\ + ##__VA_ARGS__);\ + } while(0) + } #endif //GENERIC_GRID_LAUNCH From 6c7450d414cce0f5964d4331c9f10955383ae198 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 6 Apr 2017 16:43:26 -0500 Subject: [PATCH 13/19] fixed header structure for complex data types Change-Id: I16bf19005d933f42e8c8603c5d0b2df8ea3ad04f [ROCm/hip commit: 0a07382d925e8542e4cb309b9cbcb2abe8565711] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 3 +-- projects/hip/include/hip/hcc_detail/hip_fp16.h | 2 +- projects/hip/include/hip/hcc_detail/hip_vector_types.h | 2 +- 3 files changed, 3 insertions(+), 4 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index dd742e484c..26d73a21a8 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -23,8 +23,7 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H -#include "./hip_fp16.h" -#include "./hip_vector_types.h" +#include "hip/hcc_detail/hip_vector_types.h" #if __cplusplus #define COMPLEX_ADD_OP_OVERLOAD(type) \ diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index febc1b4fce..0a861b64af 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -23,7 +23,7 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H -#include "hip/hip_runtime.h" +#include "hip/hcc_detail/hip_vector_types.h" #if __clang_major__ > 3 diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 42e1d6663c..82bd3b2d6f 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -32,7 +32,7 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif -#include "host_defines.h" +#include "hip/hcc_detail/host_defines.h" #define MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(type) \ __device__ __host__ type() {} \ From 7502166e5a42f69aab51ff81ab4f72c8cb8ef4c8 Mon Sep 17 00:00:00 2001 From: sunway513 Date: Thu, 6 Apr 2017 23:54:00 +0000 Subject: [PATCH 14/19] Improve documentation for hipModuleLaunch functions. Change-Id: I0e22621e499775740c3301347b7416d5f98c2414 [ROCm/hip commit: 599596555884903bc7670d885c2d34349190501f] --- projects/hip/include/hip/hcc_detail/hip_hcc.h | 44 ++++++++++++++++++- .../include/hip/hcc_detail/hip_runtime_api.h | 23 +++++----- 2 files changed, 54 insertions(+), 13 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_hcc.h b/projects/hip/include/hip/hcc_detail/hip_hcc.h index 645e980376..889e04eb9f 100644 --- a/projects/hip/include/hip/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hip/hcc_detail/hip_hcc.h @@ -28,6 +28,17 @@ THE SOFTWARE. #if __cplusplus #ifdef __HCC__ #include + + +/** + *------------------------------------------------------------------------------------------------- + *------------------------------------------------------------------------------------------------- + * @defgroup HCC-specific features + * @warning These APIs provide access to special features of HCC compiler and are not available through the CUDA path. + * @{ + */ + + /** * @brief Return hc::accelerator associated with the specified deviceId * @return #hipSuccess, #hipErrorInvalidDevice @@ -45,6 +56,30 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a #endif // #ifdef __HCC__ +/** + * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra + * + * @param [in[ f Kernel to launch. + * @param [in] gridDimX X grid dimension specified in work-items + * @param [in] gridDimY Y grid dimension specified in work-items + * @param [in] gridDimZ Z grid dimension specified in work-items + * @param [in] blockDimX X block dimensions specified in work-items + * @param [in] blockDimY Y grid dimension specified in work-items + * @param [in] blockDimZ Z grid dimension specified in work-items + * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules. + * @param [in] kernelParams + * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. + * @param [in] startEvent If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API. + * @param [in] stopEvent If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API. + * + * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue + + * If startNanos or stopNanos is specified, this API will record and return the start and stop timestamps for the command. The timestamps are collected on the GPU device + * and converted into ns resolution. Typically programs will specify both pointers. Collecting performance timestamps may have a small overhead (approx 1us). + * + * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage. + */ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, @@ -55,8 +90,15 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, size_t sharedMemBytes, hipStream_t hStream, void **kernelParams, - void **extra); + void **extra, + uint64_t *startNanos=nullptr, + uint64_t *stopNanos=nullptr + ); +// doxygen end HCC-specific features +/** + * @} + */ #endif // #if __cplusplus #endif // diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 0daca7a53b..f9bfb5a310 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1913,19 +1913,18 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image); /** * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra * - * @param [in[ f - * @param [in] gridDimX - * @param [in] gridDimY - * @param [in] gridDimZ - * @param [in] blockDimX - * @param [in] blockDimY - * @param [in] blockDimZ - * @param [in] sharedMemBytes - * @param [in] stream - * @param [in] kernelParams - * @param [in] extraa + * @param [in[ f Kernel to launch. + * @param [in] gridDimX X grid dimension specified as multiple of blockDimX. + * @param [in] gridDimY Y grid dimension specified as multiple of blockDimY. + * @param [in] gridDimZ Z grid dimension specified as multiple of blockDimZ. + * @param [in] blockDimX X block dimensions specified in work-items + * @param [in] blockDimY Y grid dimension specified in work-items + * @param [in] blockDimZ Z grid dimension specified in work-items + * @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The kernel can access this with HIP_DYNAMIC_SHARED. + * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules. + * @param [in] kernelParams + * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. * - * The function takes the above arguments and run the kernel in hipFunction_t f. with launch parameters specified in gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY and blockDimmZ. The amount of shared memory is specificed and can be used with HIP_DYNAMIC_SHARED. The arguemt extra is used to pass in the arguments for the kernel. * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue * * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage. From eaeecbd4619f5d19fa4bf71040a16464403149b1 Mon Sep 17 00:00:00 2001 From: sunway513 Date: Thu, 6 Apr 2017 23:55:15 +0000 Subject: [PATCH 15/19] Refactor events and add initial event option for hipHccModuleLaunchKernel - Change hipEvent_t to a class. - Move event logic inside the class. - Add _type to support Independent, StartCommand, StopCommand events. StartCommand returns start timestamp from events. Change-Id: I4ddd694f2645a3ff7170c9111dc1d3e39931ca21 [ROCm/hip commit: cfa3155082b6ab3eb91b78cbed3e269e9a846ce4] --- projects/hip/src/hip_event.cpp | 60 +++++++++++++++++++++++++---- projects/hip/src/hip_hcc.cpp | 17 -------- projects/hip/src/hip_hcc_internal.h | 31 +++++++++++---- projects/hip/src/hip_module.cpp | 32 +++++++++++---- 4 files changed, 99 insertions(+), 41 deletions(-) diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index d44f201db5..61ac5cd3ab 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -30,6 +30,54 @@ THE SOFTWARE. //--- +ihipEvent_t::ihipEvent_t(unsigned flags) +{ + _state = hipEventStatusCreated; + _stream = NULL; + _flags = flags; + _timestamp = 0; + _type = hipEventTypeIndependent; +}; + + + +// Attach to an existing completion future: +void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, ihipEventType_t eventType) +{ + _state = hipEventStatusRecording; + _marker = *cf; + _type = eventType; +} + + + +void ihipEvent_t::setTimestamp() +{ + if (_state == hipEventStatusRecorded) { + // already recorded, done: + return; + } else { + // TODO - use completion-future functions to obtain ticks and timestamps: + hsa_signal_t *sig = static_cast (_marker.get_native_handle()); + if (sig) { + if (hsa_signal_load_acquire(*sig) == 0) { + + if ((_type == hipEventTypeIndependent) || (_type == hipEventTypeStopCommand)) { + _timestamp = _marker.get_end_tick(); + } else if (_type == hipEventTypeStartCommand) { + _timestamp = _marker.get_begin_tick(); + } else { + assert(0); // TODO - move to debug assert + _timestamp = 0; + } + + _state = hipEventStatusRecorded; + } + } + } +} + + hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) { hipError_t e = hipSuccess; @@ -37,12 +85,8 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) // TODO-IPC - support hipEventInterprocess. unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming; if ((flags & ~supportedFlags) == 0) { - ihipEvent_t *eh = new ihipEvent_t(); + ihipEvent_t *eh = new ihipEvent_t(flags); - eh->_state = hipEventStatusCreated; - eh->_stream = NULL; - eh->_flags = flags; - eh->_timestamp = 0; *event = eh; } else { e = hipErrorInvalidValue; @@ -141,8 +185,8 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) ihipEvent_t *start_eh = start; ihipEvent_t *stop_eh = stop; - ihipSetTs(start); - ihipSetTs(stop); + start->setTimestamp(); + stop->setTimestamp(); hipError_t status = hipSuccess; *ms = 0.0f; @@ -151,7 +195,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) if ((start_eh->_state == hipEventStatusRecorded) && (stop_eh->_state == hipEventStatusRecorded)) { // Common case, we have good information for both events. - int64_t tickDiff = (stop_eh->_timestamp - start_eh->_timestamp); + int64_t tickDiff = (stop_eh->timestamp() - start_eh->timestamp()); uint64_t freqHz; hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 374840f91f..35a3e11e71 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -1641,23 +1641,6 @@ const char *ihipErrorString(hipError_t hip_error) }; -void ihipSetTs(hipEvent_t e) -{ - ihipEvent_t *eh = e; - if (eh->_state == hipEventStatusRecorded) { - // already recorded, done: - return; - } else { - // TODO - use completion-future functions to obtain ticks and timestamps: - hsa_signal_t *sig = static_cast (eh->_marker.get_native_handle()); - if (sig) { - if (hsa_signal_load_acquire(*sig) == 0) { - eh->_timestamp = eh->_marker.get_end_tick(); - eh->_state = hipEventStatusRecorded; - } - } - } -} // Returns true if copyEngineCtx can see the memory allocated on dstCtx and srcCtx. diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 4b960e2820..459ea3ba2c 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -584,22 +584,39 @@ private: // Data //---- // Internal event structure: enum hipEventStatus_t { - hipEventStatusUnitialized = 0, // event is unutilized, must be "Created" before use. - hipEventStatusCreated = 1, - hipEventStatusRecording = 2, // event has been enqueued to record something. - hipEventStatusRecorded = 3, // event has been recorded - timestamps are valid. + hipEventStatusUnitialized = 0, // event is unutilized, must be "Created" before use. + hipEventStatusCreated = 1, + hipEventStatusRecording = 2, // event has been enqueued to record something. + hipEventStatusRecorded = 3, // event has been recorded - timestamps are valid. } ; +// TODO - rename to ihip type of some kind +enum ihipEventType_t { + hipEventTypeIndependent, + hipEventTypeStartCommand, + hipEventTypeStopCommand, +}; // internal hip event structure. -struct ihipEvent_t { - hipEventStatus_t _state; +class ihipEvent_t { +public: + ihipEvent_t(unsigned flags); + void attachToCompletionFuture(const hc::completion_future *cf, ihipEventType_t eventType); + void setTimestamp(); + uint64_t timestamp() const { return _timestamp; } ; + +public: + hipEventStatus_t _state; hipStream_t _stream; // Stream where the event is recorded, or NULL if all streams. unsigned _flags; hc::completion_future _marker; + +private: + ihipEventType_t _type; uint64_t _timestamp; // store timestamp, may be set on host or by marker. +friend hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream); } ; @@ -822,8 +839,6 @@ extern hipError_t ihipDeviceSetState(); extern ihipDevice_t *ihipGetDevice(int); ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex); -extern void ihipSetTs(hipEvent_t e); - hipStream_t ihipSyncAndResolveStream(hipStream_t); diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 67bba5f935..c8555672c3 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -364,10 +364,11 @@ hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, hipError_t ihipModuleLaunchKernel(hipFunction_t f, - uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, - uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, - size_t sharedMemBytes, hipStream_t hStream, - void **kernelParams, void **extra) + uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, + size_t sharedMemBytes, hipStream_t hStream, + void **kernelParams, void **extra, + hipEvent_t *startEvent, hipEvent_t *stopEvent) { auto ctx = ihipGetTlsDefaultCtx(); @@ -446,7 +447,20 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); }; - lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, nullptr/*completion_future*/); + + hc::completion_future cf; + + lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, + (startEvent || stopEvent) ? &cf : nullptr); + + + if (startEvent) { + (*startEvent)->attachToCompletionFuture(&cf, hipEventTypeStartCommand); + } + if (stopEvent) { + (*stopEvent)->attachToCompletionFuture (&cf, hipEventTypeStopCommand); + } + if(kernelParams != NULL){ free(config[1]); @@ -470,7 +484,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, return ihipLogStatus(ihipModuleLaunchKernel(f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY, blockDimZ, - sharedMemBytes, hStream, kernelParams, extra)); + sharedMemBytes, hStream, kernelParams, extra, + nullptr, nullptr)); } @@ -478,7 +493,8 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, - void **kernelParams, void **extra) + void **kernelParams, void **extra, + hipEvent_t *startEvent, hipEvent_t *stopEvent) { HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, @@ -486,7 +502,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, kernelParams, extra); return ihipLogStatus(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, - sharedMemBytes, hStream, kernelParams, extra)); + sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent)); } hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, From cbfc58660ea9f04184fff4d940013848cf0cfc7e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 6 Apr 2017 21:00:00 -0500 Subject: [PATCH 16/19] Finish adding start/stop event support to hipHccModuleLaunchKernel. Change interface to use hipEvent_t rather than hipEvent_t* Change-Id: I259062dc087a13d51dc27f84e1e8861f332a104d [ROCm/hip commit: 6656d33f75bbc13c48e6e24b02ec43dc43af1fac] --- projects/hip/include/hip/hcc_detail/hip_hcc.h | 11 +++++------ projects/hip/src/hip_hcc_internal.h | 1 + projects/hip/src/hip_module.cpp | 8 ++++---- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_hcc.h b/projects/hip/include/hip/hcc_detail/hip_hcc.h index 889e04eb9f..fc04917931 100644 --- a/projects/hip/include/hip/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hip/hcc_detail/hip_hcc.h @@ -70,15 +70,14 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a * @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules. * @param [in] kernelParams * @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel. - * @param [in] startEvent If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API. + * @param [in] startEvent If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API. * @param [in] stopEvent If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API. * * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue - - * If startNanos or stopNanos is specified, this API will record and return the start and stop timestamps for the command. The timestamps are collected on the GPU device - * and converted into ns resolution. Typically programs will specify both pointers. Collecting performance timestamps may have a small overhead (approx 1us). * * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage. + + * HIP/ROCm actually updates the start event when the associated kernel completes. */ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, @@ -91,8 +90,8 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, hipStream_t hStream, void **kernelParams, void **extra, - uint64_t *startNanos=nullptr, - uint64_t *stopNanos=nullptr + hipEvent_t startEvent=nullptr, + hipEvent_t stopEvent=nullptr ); // doxygen end HCC-specific features diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 459ea3ba2c..9c17c6e98c 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -604,6 +604,7 @@ public: void attachToCompletionFuture(const hc::completion_future *cf, ihipEventType_t eventType); void setTimestamp(); uint64_t timestamp() const { return _timestamp; } ; + ihipEventType_t type() const { return _type; }; public: hipEventStatus_t _state; diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index c8555672c3..b359e7a63c 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -368,7 +368,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra, - hipEvent_t *startEvent, hipEvent_t *stopEvent) + hipEvent_t startEvent, hipEvent_t stopEvent) { auto ctx = ihipGetTlsDefaultCtx(); @@ -455,10 +455,10 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, if (startEvent) { - (*startEvent)->attachToCompletionFuture(&cf, hipEventTypeStartCommand); + startEvent->attachToCompletionFuture(&cf, hipEventTypeStartCommand); } if (stopEvent) { - (*stopEvent)->attachToCompletionFuture (&cf, hipEventTypeStopCommand); + stopEvent->attachToCompletionFuture (&cf, hipEventTypeStopCommand); } @@ -494,7 +494,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra, - hipEvent_t *startEvent, hipEvent_t *stopEvent) + hipEvent_t startEvent, hipEvent_t stopEvent) { HIP_INIT_API(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, From 6825ff455d1d458bff6da8e0af714ee6d48ebb81 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 7 Apr 2017 14:51:54 +0530 Subject: [PATCH 17/19] Fix build issues in hipCommander sample - Remove -stdlib=libstdc++ from Makefile - Removed deleted HIP header file fom includes Change-Id: Ia189396bee19fc52b679259df56c6c6e2bafb6fe [ROCm/hip commit: 6e565d8469e9fe4a3da95ee60155dfe5ca85e95b] --- projects/hip/samples/1_Utils/hipCommander/Makefile | 3 --- projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp | 1 - 2 files changed, 4 deletions(-) diff --git a/projects/hip/samples/1_Utils/hipCommander/Makefile b/projects/hip/samples/1_Utils/hipCommander/Makefile index e770c636a4..a411763b7f 100644 --- a/projects/hip/samples/1_Utils/hipCommander/Makefile +++ b/projects/hip/samples/1_Utils/hipCommander/Makefile @@ -10,9 +10,6 @@ OPT=-O3 CXXFLAGS = $(OPT) --std=c++11 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) -ifeq (${HIP_PLATFORM}, hcc) - CXXFLAGS += " -stdlib=libc++" -endif CODE_OBJECTS=nullkernel.hsaco diff --git a/projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp index 0add1ce3e3..4b93180b18 100644 --- a/projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/projects/hip/samples/1_Utils/hipCommander/hipCommander.cpp @@ -11,7 +11,6 @@ #include #include #include -#include #endif #include From acbde105c64bc494c28381c82d2db41951c4fb5a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 7 Apr 2017 15:24:10 +0530 Subject: [PATCH 18/19] Fix build issues with bit_extract sample Change-Id: I628b3c83a16f7adf0ab8ca60aecde8c073c34fd9 [ROCm/hip commit: 935e3cd64975274f322d0b5c7064e6b1fc6f49b9] --- projects/hip/samples/0_Intro/bit_extract/Makefile | 4 ---- 1 file changed, 4 deletions(-) diff --git a/projects/hip/samples/0_Intro/bit_extract/Makefile b/projects/hip/samples/0_Intro/bit_extract/Makefile index 78f6a2faa8..08bca6e642 100644 --- a/projects/hip/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip/samples/0_Intro/bit_extract/Makefile @@ -11,10 +11,6 @@ HIPCC=$(HIP_PATH)/bin/hipcc ifeq (${HIP_PLATFORM}, nvcc) HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 endif -ifeq (${HIP_PLATFORM}, hcc) - HIPCC_FLAGS = -stdlib=libc++ -endif - EXE=bit_extract From 3e976c9aa09c1ef5d01881ca72a1edf13cb94c79 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 7 Apr 2017 15:38:56 +0530 Subject: [PATCH 19/19] Updated table of contents in markdown documentation Change-Id: I7347a06f57f9927ca3fcc5590a6c8200bc1bb1f5 [ROCm/hip commit: 16215ea9b2718f3d2a98f5d208a4906b0a878017] --- projects/hip/docs/markdown/hip_bugs.md | 6 +-- projects/hip/docs/markdown/hip_faq.md | 9 ++-- .../hip/docs/markdown/hip_kernel_language.md | 1 + .../hip/docs/markdown/hip_porting_guide.md | 9 ++-- projects/hip/docs/markdown/hip_profiling.md | 46 +++++++++++-------- 5 files changed, 41 insertions(+), 30 deletions(-) diff --git a/projects/hip/docs/markdown/hip_bugs.md b/projects/hip/docs/markdown/hip_bugs.md index 14f2935f17..73133843bc 100644 --- a/projects/hip/docs/markdown/hip_bugs.md +++ b/projects/hip/docs/markdown/hip_bugs.md @@ -1,10 +1,10 @@ -# HIP Bugs +# HIP Bugs -- [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**](#errors-related-to-undefined-reference-to-hclaunchkernel__grid_launch_parm) -- [Application hangs after a hipLaunchKernel call](#what-if-i-see-application-hangs-after-a-hiplaunchkernel-call) +- [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**`](#errors-related-to-undefined-reference-to-__hclaunchkernel____grid_launch_parm) - [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method) +- [Errors related to `no matching constructor`](#errors-related-to-no-matching-constructor) - [HIP is more restrictive in enforcing restrictions](#hip-is-more-restrictive-in-enforcing-restrictions) diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md index 8ccb458103..e316d449ef 100644 --- a/projects/hip/docs/markdown/hip_faq.md +++ b/projects/hip/docs/markdown/hip_faq.md @@ -4,7 +4,7 @@ - [What APIs and features does HIP support?](#what-apis-and-features-does-hip-support) - [What is not supported?](#what-is-not-supported) - * [Run-time features](#run-time-features) + * [Runtime/Driver API features](#runtimedriver-api-features) * [Kernel language features](#kernel-language-features) - [Is HIP a drop-in replacement for CUDA?](#is-hip-a-drop-in-replacement-for-cuda) - [What specific version of CUDA does HIP support?](#what-specific-version-of-cuda-does-hip-support) @@ -23,10 +23,11 @@ - [On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ?](#on-hcc-can-i-link-hip-code-with-host-code-compiled-with-another-compiler-such-as-gcc-icc-or-clang-) - [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do) - [Can I install both CUDA SDK and HCC on same machine?](#can-i-install-both-cuda-sdk-and-hcc-on-same-machine) +- [On CUDA, can I mix CUDA code with HIP code?](#on-cuda-can-i-mix-cuda-code-with-hip-code) +- [On HCC, can I use HC functionality with HIP?](#on-hcc-can-i-use-hc-functionality-with-hip) - [How do I trace HIP application flow?](#how-do-i-trace-hip-application-flow) - * [Using CodeXL markers for HIP Functions](#using-codexl-markers-for-hip-functions) - * [Using HIP_TRACE_API](#using-hip_trace_api) -- [How do I enable HIP Generic Grid Launch option?](#how-do-i-enable-hip-generic-grid-launch-option) +- [What if HIP generates error of "symbol multiply defined!" only on AMD machine?](#what-if-hip-generates-error-of-symbol-multiply-defined-only-on-amd-machine) +- [How do I disable HIP Generic Grid Launch option?](#how-do-i-disable-hip-generic-grid-launch-option) diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index 0c7f3c8d25..3cb7b17a0c 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -44,6 +44,7 @@ - [Pragma Unroll](#pragma-unroll) - [In-Line Assembly](#in-line-assembly) - [C++ Support](#c-support) +- [Kernel Compilation](#kernel-compilation) diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 9f20d12423..72f6384f6d 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -21,6 +21,7 @@ and provides practical suggestions on how to port CUDA code and work through com * [Device-Architecture Properties](#device-architecture-properties) * [Table of Architecture Properties](#table-of-architecture-properties) - [Finding HIP](#finding-hip) +- [hipLaunchKernel](#hiplaunchkernel) - [Compiler Options](#compiler-options) - [Linking Issues](#linking-issues) * [Linking With hipcc](#linking-with-hipcc) @@ -31,9 +32,11 @@ and provides practical suggestions on how to port CUDA code and work through com * [Using a Standard C++ Compiler](#using-a-standard-c-compiler) + [cuda.h](#cudah) * [Choosing HIP File Extensions](#choosing-hip-file-extensions) - * [Workarounds](#workarounds) - + [warpSize](#warpsize) - + [Textures and Cache Control](#textures-and-cache-control) +- [Workarounds](#workarounds) + * [warpSize](#warpsize) +- [memcpyToSymbol](#memcpytosymbol) +- [threadfence_system](#threadfence_system) + * [Textures and Cache Control](#textures-and-cache-control) - [More Tips](#more-tips) * [HIPTRACE Mode](#hiptrace-mode) * [Environment Variables](#environment-variables) diff --git a/projects/hip/docs/markdown/hip_profiling.md b/projects/hip/docs/markdown/hip_profiling.md index 463c9c13b3..6e5cde700d 100644 --- a/projects/hip/docs/markdown/hip_profiling.md +++ b/projects/hip/docs/markdown/hip_profiling.md @@ -4,26 +4,32 @@ This section describes the profiling and debugging capabilities that HIP provide Profiling information can viewed in the CodeXL visualization tool or printed directly to stderr as the application runs. This document starts with some of the general capabilities of CodeXL and then describes some of the additional HIP marker and debug features. - * [CodeXL Profiling](#codexl-profiling) - * [Collecting and Viewing Traces](#collecting-and-viewing-traces) - * [Using rocm-profiler timestamp profiling](#using-rocm-profiler-timestamp-profiling) - * [Using rocm-profiler performance counter collection:](#using-rocm-profiler-performance-counter-collection) - * [Using CodeXL to view profiling results:](#using-codexl-to-view-profiling-results) - * [More information on CodeXL](#more-information-on-codexl) - * [HIP Markers](#hip-markers) - * [Profiling HIP APIs](#profiling-hip-apis) - * [Adding markers to applications](#adding-markers-to-applications) - * [Additional HIP Profiling Features](#additional-hip-profiling-features) - * [Demangling C Kernel Names](#demangling-c-kernel-names) - * [Controlling when profiling starts and ends](#controlling-when-profiling-starts-and-ends) - * [Reducing timeline trace output file size](#reducing-timeline-trace-output-file-size) - * [How to enable profiling at HIP build time](#how-to-enable-profiling-at-hip-build-time) - * [Tracing and Debug](#tracing-and-debug) - * [Tracing HIP APIs](#tracing-hip-apis) - * [Color](#color) - * [Using HIP_DB](#using-hip_db) - * [Using ltrace](#using-ltrace) - * [Chicken bits](#chicken-bits) + + +- [CodeXL Profiling](#codexl-profiling) + * [Collecting and Viewing Traces](#collecting-and-viewing-traces) + + [Using rocm-profiler timestamp profiling](#using-rocm-profiler-timestamp-profiling) + + [Using rocm-profiler performance counter collection:](#using-rocm-profiler-performance-counter-collection) + + [Using CodeXL to view profiling results:](#using-codexl-to-view-profiling-results) + + [More information on CodeXL](#more-information-on-codexl) + * [HIP Markers](#hip-markers) + + [Profiling HIP APIs](#profiling-hip-apis) + + [Adding markers to applications](#adding-markers-to-applications) + * [Additional HIP Profiling Features](#additional-hip-profiling-features) + + [Demangling C++ Kernel Names](#demangling-c-kernel-names) + + [Controlling when profiling starts and ends](#controlling-when-profiling-starts-and-ends) + + [Reducing timeline trace output file size](#reducing-timeline-trace-output-file-size) + + [How to enable profiling at HIP build time](#how-to-enable-profiling-at-hip-build-time) +- [Tracing and Debug](#tracing-and-debug) + * [Tracing HIP APIs](#tracing-hip-apis) + + [Color](#color) + * [Using HIP_DB](#using-hip_db) + * [Using ltrace](#using-ltrace) + * [Chicken bits](#chicken-bits) + * [Debugging HIP Applications](#debugging-hip-applications) + * [General Debugging Tips](#general-debugging-tips) + + ## CodeXL Profiling