From 1003128925d5e4dd1870763f26e97316c66d3b01 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 20 May 2020 10:42:38 -0400 Subject: [PATCH 1/9] Don't search ROCclr_DIR for prof_protocol.h This points to the cmake directory where the find module was found, not a prefix for where it was found. Based on the search below looking in roctracer, searching in ROCclr for the header doesn't make much sense. The header should be either provided by ROCclr xor roctracer. Having it possibly be provided by two different dependencies is confusing, and a potential source of version mismatch problems. Change-Id: Ic2f6ec03f9a7b86225cf7e5c43f39a1360318a34 --- hipamd/rocclr/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/hipamd/rocclr/CMakeLists.txt b/hipamd/rocclr/CMakeLists.txt index 9d05de586d..a6f918705a 100644 --- a/hipamd/rocclr/CMakeLists.txt +++ b/hipamd/rocclr/CMakeLists.txt @@ -58,7 +58,6 @@ if (NOT ROCclr_FOUND) include(${LIBROCclr_STATIC_DIR}/amdrocclr_staticTargets.cmake) endif() -set(PROF_API_HEADER_PATH ${ROCclr_DIR}/platform) ############################# # Profiling API support ############################# From 3a30b9eb8dc75502fd807bdac80cb422ec1c0288 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 15 May 2020 18:08:26 -0400 Subject: [PATCH 2/9] Add support for hipExtStreamCreateWithCUMask API Change-Id: I369d0eaca493821c4badc6b18ac02daa2fddc95f --- .../include/hip/hcc_detail/hip_runtime_api.h | 22 ++++++++++++++ hipamd/rocclr/hip_hcc.def.in | 1 + hipamd/rocclr/hip_hcc.map.in | 1 + hipamd/rocclr/hip_internal.hpp | 4 ++- hipamd/rocclr/hip_stream.cpp | 29 +++++++++++++++---- 5 files changed, 51 insertions(+), 6 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 744885923f..d908f72a04 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -849,6 +849,28 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags); hipError_t hipStreamGetPriority(hipStream_t stream, int* priority); +/** + * @brief Create an asynchronous stream with the specified CU mask. + * + * @param[in, out] stream Pointer to new stream + * @param[in ] cuMaskSize Size of CU mask bit array passed in. + * @param[in ] cuMask Bit-vector representing the CU mask. Each active bit represents using one CU. + * The first 32 bits represent the first 32 CUs, and so on. If its size is greater than physical + * CU number (i.e., multiProcessorCount member of hipDeviceProp_t), the extra elements are ignored. + * It is user's responsibility to make sure the input is meaningful. + * @return #hipSuccess, #hipErrorInvalidHandle, #hipErrorInvalidValue + * + * Create a new asynchronous stream with the specified CU mask. @p stream returns an opaque handle + * that can be used to reference the newly created stream in subsequent hipStream* commands. The + * stream is allocated on the heap and will remain allocated even if the handle goes out-of-scope. + * To release the memory used by the stream, application must call hipStreamDestroy. + * + * + * @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy + */ +hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, const uint32_t* cuMask); + + /** * Stream CallBack struct */ diff --git a/hipamd/rocclr/hip_hcc.def.in b/hipamd/rocclr/hip_hcc.def.in index 579608e685..53bdfb3065 100755 --- a/hipamd/rocclr/hip_hcc.def.in +++ b/hipamd/rocclr/hip_hcc.def.in @@ -252,3 +252,4 @@ hipTexObjectDestroy hipTexObjectGetResourceDesc hipTexObjectGetResourceViewDesc hipTexObjectGetTextureDesc +hipExtStreamCreateWithCUMask diff --git a/hipamd/rocclr/hip_hcc.map.in b/hipamd/rocclr/hip_hcc.map.in index 19da8a6991..02a704bf88 100755 --- a/hipamd/rocclr/hip_hcc.map.in +++ b/hipamd/rocclr/hip_hcc.map.in @@ -258,6 +258,7 @@ global: hipInitActivityCallback*; hipEnableActivityCallback*; hipGetCmdName*; + hipExtStreamCreateWithCUMask; }; local: *; diff --git a/hipamd/rocclr/hip_internal.hpp b/hipamd/rocclr/hip_internal.hpp index 4cc0dadd8a..492842c32b 100755 --- a/hipamd/rocclr/hip_internal.hpp +++ b/hipamd/rocclr/hip_internal.hpp @@ -87,9 +87,11 @@ namespace hip { amd::CommandQueue::Priority priority_; unsigned int flags_; bool null_; + const std::vector cuMask_; public: - Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false); + Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false, + const std::vector& cuMask = {}); /// Creates the hip stream object, including AMD host queue bool Create(); diff --git a/hipamd/rocclr/hip_stream.cpp b/hipamd/rocclr/hip_stream.cpp index 3bd7d343f7..23b89e71ca 100644 --- a/hipamd/rocclr/hip_stream.cpp +++ b/hipamd/rocclr/hip_stream.cpp @@ -44,15 +44,15 @@ namespace hip { // ================================================================================================ Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p, - unsigned int f, bool null_stream) + unsigned int f, bool null_stream, const std::vector& cuMask) : queue_(nullptr), lock_("Stream Callback lock"), device_(dev), - priority_(p), flags_(f), null_(null_stream) {} + priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {} // ================================================================================================ bool Stream::Create() { cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; queue_ = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties, - amd::CommandQueue::RealTimeDisabled, priority_); + amd::CommandQueue::RealTimeDisabled, priority_, cuMask_); // Create a host queue bool result = (queue_ != nullptr) ? queue_->create() : false; // Insert just created stream into the list of the blocking queues @@ -160,8 +160,9 @@ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, // ================================================================================================ static hipError_t ihipStreamCreate(hipStream_t* stream, - unsigned int flags, amd::CommandQueue::Priority priority) { - hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags); + unsigned int flags, amd::CommandQueue::Priority priority, + const std::vector& cuMask = {}) { + hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask); if (hStream == nullptr) { return hipErrorOutOfMemory; @@ -311,3 +312,21 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback HIP_RETURN(hipSuccess); } + +// ================================================================================================ +hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, + const uint32_t* cuMask) { + HIP_INIT_API(hipExtStreamCreateWithCUMask, stream, cuMaskSize, cuMask); + + if (stream == nullptr) { + HIP_RETURN(hipErrorInvalidHandle); + } + if (cuMaskSize == 0 || cuMask == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + const std::vector cuMaskv(cuMask, cuMask + cuMaskSize); + + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv)); +} + From 51b6c1856c89098536f292b02fa90bbdcaaaaedf Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Mon, 18 May 2020 16:17:34 -0400 Subject: [PATCH 3/9] Fix square build failure with static lib in Jenkin When libamdhip64_static.a is built by Jenkin, sample square cannot been built successfully because libamdhip64_static.a is archiveved in thin mode. Thus in the patch it will be archiveved in full mode. Meanwhile libamdhip64_static_temp.a will be useless and thus removed. Change-Id: Ifd3882598ef0dc5e7af8db0e389e786025ceb455 --- hipamd/rocclr/CMakeLists.txt | 21 +++++---------------- 1 file changed, 5 insertions(+), 16 deletions(-) diff --git a/hipamd/rocclr/CMakeLists.txt b/hipamd/rocclr/CMakeLists.txt index a6f918705a..2edcdabb55 100644 --- a/hipamd/rocclr/CMakeLists.txt +++ b/hipamd/rocclr/CMakeLists.txt @@ -203,17 +203,6 @@ set_target_properties(hip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR}) set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR}) set_target_properties(amdhip64_static PROPERTIES PUBLIC_HEADER ${PROF_API_STR}) - - -# We expect amdhip64_static to contain objects of rocclr and hip. But linker -# let amdhip64_static contain objects of hip only. So we will use a -# a custom amdhip64_static_combiner to combine objects of vid and hip into -# amdhip64_static. To avoid amdhip64_static contains itself, -# amdhip64_static_temp is created internally. -add_library(amdhip64_static_temp STATIC - $ - ) - add_library(host INTERFACE) target_link_libraries(host INTERFACE amdhip64) add_library(device INTERFACE) @@ -225,19 +214,19 @@ target_link_libraries(device INTERFACE host) # filename. target_link_libraries(amdhip64 PRIVATE amdrocclr_static Threads::Threads dl) target_link_libraries(amdhip64_static PRIVATE Threads::Threads dl) -target_link_libraries(amdhip64_static_temp PRIVATE Threads::Threads dl) # combine objects of vid and hip into amdhip64_static add_custom_target( amdhip64_static_combiner ALL - COMMAND rm -f $ # Must remove old one, otherwise the new one will contain obsolete stuff - COMMAND ${CMAKE_AR} -rcsT $ $ $ - DEPENDS amdhip64_static amdhip64_static_temp amdrocclr_static # To make sure this is the last step + COMMAND rm -rf static_lib_temp && mkdir static_lib_temp && cd static_lib_temp # Create temp folder to contain *.o + COMMAND ${CMAKE_AR} -x $ # Extract *.o from amdrocclr_static + COMMAND ${CMAKE_AR} -rcs $ *.o # Append *.o to amdhip64_static + COMMAND cd .. && rm -rf static_lib_temp # Remove temp folder + DEPENDS amdhip64_static amdrocclr_static # To make sure this is the last step COMMENT "Combining static libs into amdhip64_static" ) - INSTALL(PROGRAMS $ DESTINATION lib COMPONENT MAIN) INSTALL(PROGRAMS $ DESTINATION lib COMPONENT MAIN) INSTALL(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhip_hcc.so )" DESTINATION lib COMPONENT MAIN) From 355661b5da583fb66a43ae7696fdc69884eb49b7 Mon Sep 17 00:00:00 2001 From: Vlad Sytchenko Date: Fri, 22 May 2020 14:13:50 -0400 Subject: [PATCH 4/9] Reenable texture reference tests Change-Id: I77024476cff77951d61dc48f7e30094d6b47266c --- hipamd/tests/src/texture/hipBindTex2DPitch.cpp | 2 +- hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp | 2 +- .../src/texture/hipNormalizedFloatValueTex.cpp | 14 +------------- hipamd/tests/src/texture/hipTextureRef2D.cpp | 2 +- .../tests/src/texture/simpleTexture2DLayered.cpp | 2 +- hipamd/tests/src/texture/simpleTexture3D.cpp | 2 +- 6 files changed, 6 insertions(+), 18 deletions(-) diff --git a/hipamd/tests/src/texture/hipBindTex2DPitch.cpp b/hipamd/tests/src/texture/hipBindTex2DPitch.cpp index 2fd3f1228d..3be39fb911 100644 --- a/hipamd/tests/src/texture/hipBindTex2DPitch.cpp +++ b/hipamd/tests/src/texture/hipBindTex2DPitch.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /*HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ diff --git a/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp b/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp index eeaf42129f..546e9775e8 100644 --- a/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp +++ b/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ diff --git a/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp b/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp index 95b7c1d879..f365240543 100644 --- a/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/hipamd/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc hcc rocclr + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc hcc * TEST: %t * HIT_END */ @@ -42,24 +42,12 @@ static float getNormalizedValue(const float value, return value; } -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texuc; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texs; -#if __HIP__ -__hip_pinned_shadow__ -#endif texture texus; diff --git a/hipamd/tests/src/texture/hipTextureRef2D.cpp b/hipamd/tests/src/texture/hipTextureRef2D.cpp index 4f00260998..e39c36638c 100644 --- a/hipamd/tests/src/texture/hipTextureRef2D.cpp +++ b/hipamd/tests/src/texture/hipTextureRef2D.cpp @@ -1,5 +1,5 @@ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ diff --git a/hipamd/tests/src/texture/simpleTexture2DLayered.cpp b/hipamd/tests/src/texture/simpleTexture2DLayered.cpp index 1424976716..d8db1308e3 100644 --- a/hipamd/tests/src/texture/simpleTexture2DLayered.cpp +++ b/hipamd/tests/src/texture/simpleTexture2DLayered.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr + * BUILD: %t %s ../test_common.cpp * TEST: %t * HIT_END */ diff --git a/hipamd/tests/src/texture/simpleTexture3D.cpp b/hipamd/tests/src/texture/simpleTexture3D.cpp index bd342df1c3..ca6695f619 100644 --- a/hipamd/tests/src/texture/simpleTexture3D.cpp +++ b/hipamd/tests/src/texture/simpleTexture3D.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc rocclr + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ From a98920d9a3df87f8bfd0bce3df0721bdf5e39e07 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Tue, 12 May 2020 10:10:58 -0400 Subject: [PATCH 5/9] Remove ROCclr search hacks find_package should now be the only way to import ROCclr. Also update the build example comment. The build scripts used 2 custom variables to manually specify the build and source directories for where to find VDI. Once renamed to ROCclr, these conflicted with the variables automatically set by find_package(ROCclr). These hacks tried to satisfy this intermediate step to try satisfying commit ordering problems to get through PSDB. The INSTALL.md documentation should also be updated, but it's completely missing any mention of ROCclr now, and still gives directions for hcc. Change-Id: I6fc94b6cb36241a9d4f22d24e49523367f803461 --- hipamd/CMakeLists.txt | 2 +- hipamd/rocclr/CMakeLists.txt | 36 +----------------------------------- 2 files changed, 2 insertions(+), 36 deletions(-) diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index 16c3f11edf..8278daad2c 100755 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.4.3) project(hip) # sample command for hip-rocclr, you'll need to have rocclr installed # cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr .. -# cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr -DROCclr_DIR=/extra/lmoriche/hip-rocclr/rocclr -DOPENCL_DIR=/extra/lmoriche/clients/lmoriche_opencl_dev2/drivers/opencl/api/opencl -DLIBROCclr_STATIC_DIR=/extra/lmoriche/hip-rocclr/build/rocclr .. +# cmake -DHIP_COMPILER=clang -DHIP_PLATFORM=rocclr -DOPENCL_DIR=/path/to/opencl/api/opencl -DCMAKE_PREFIX_PATH=/path/to/rocclr/build/or/install/directory .. set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake") diff --git a/hipamd/rocclr/CMakeLists.txt b/hipamd/rocclr/CMakeLists.txt index 2edcdabb55..37135af81d 100644 --- a/hipamd/rocclr/CMakeLists.txt +++ b/hipamd/rocclr/CMakeLists.txt @@ -27,37 +27,11 @@ endif() set(USE_PROF_API "1") -# FIXME: Make this required and remove the legacy handling below -set(save_rocclr_dir ${ROCclr_DIR}) -set(save_rocclr_static_dir ${LIBROCclr_STATIC_DIR}) - -find_package(ROCclr CONFIG +find_package(ROCclr REQUIRED CONFIG PATHS /opt/rocm /opt/rocm/rocclr) -if (NOT ROCclr_FOUND) - if(NOT DEFINED LIBROCclr_STATIC_DIR) - find_path(LIBROCclr_STATIC_DIR - NAMES libamdrocclr_static.a - PATHS /opt/rocm/rocclr - PATH_SUFFIXES lib) - else() - set(LIBROCclr_STATIC_DIR ${save_rocclr_static_dir}) - endif() - - if(NOT DEFINED ROCclr_DIR) - find_path(ROCclr_DIR - NAMES top.hpp - PATH_SUFFIXES include - PATHS /opt/rocm/rocclr) - else() - set(ROCclr_DIR ${save_rocclr_dir}) - endif() - message("Found Static rocclr lib:${LIBROCclr_STATIC_DIR} and rocclr includes: ${ROCclr_DIR}") - include(${LIBROCclr_STATIC_DIR}/amdrocclr_staticTargets.cmake) -endif() - ############################# # Profiling API support ############################# @@ -131,14 +105,6 @@ target_include_directories(hip64 ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/amdocl ${PROJECT_SOURCE_DIR}/include/hip/hcc_detail/elfio - - # FIXME: Remove ROCclr_DIr explicit references - ${ROCclr_DIR} - ${ROCclr_DIR}/include - ${ROCclr_DIR}/compiler/lib - ${ROCclr_DIR}/compiler/lib/include - ${ROCclr_DIR}/elf/utils/common - ${ROCclr_DIR}/elf/utils/libelf ${ROCR_INCLUDES} $ $) From f4e6dec3acd49c31a1fe454bedb3170e83916470 Mon Sep 17 00:00:00 2001 From: Mahesha Shivamallappa Date: Wed, 13 May 2020 14:37:49 +0530 Subject: [PATCH 6/9] Add support for cooperative group type - thread_block Change-Id: If3770b6d6718a638b70f527ae2533d9ef3267ff4 --- .../hip/hcc_detail/hip_cooperative_groups.h | 59 +++++++ .../hip_cooperative_groups_helper.h | 40 ++++- hipamd/tests/src/cg/hipCGThreadBlockType.cpp | 164 ++++++++++++++++++ 3 files changed, 262 insertions(+), 1 deletion(-) create mode 100644 hipamd/tests/src/cg/hipCGThreadBlockType.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_cooperative_groups.h b/hipamd/include/hip/hcc_detail/hip_cooperative_groups.h index 27ce887723..73f6778a1c 100644 --- a/hipamd/include/hip/hcc_detail/hip_cooperative_groups.h +++ b/hipamd/include/hip/hcc_detail/hip_cooperative_groups.h @@ -166,6 +166,55 @@ this_grid() { return grid_group(internal::grid::size()); } +/** \brief The workgroup (thread-block in CUDA terminology) cooperative group + * type + * + * \details Represents an intra-workgroup cooperative group type where the + * participating threads within the group are exctly the same threads + * which are participated in the currently executing `workgroup` + */ +class thread_block : public thread_group { + // Only these friend functions are allowed to construct an object of this + // class and access its resources + friend __CG_QUALIFIER__ thread_block this_thread_block(); + + protected: + // Construct a workgroup thread group (through the API this_thread_block()) + explicit __CG_QUALIFIER__ thread_block(uint32_t size) + : thread_group(internal::cg_workgroup, size) { } + + public: + // 3-dimensional block index within the grid + __CG_QUALIFIER__ dim3 group_index() { + return internal::workgroup::group_index(); + } + // 3-dimensional thread index within the block + __CG_QUALIFIER__ dim3 thread_index() { + return internal::workgroup::thread_index(); + } + __CG_QUALIFIER__ uint32_t thread_rank() const { + return internal::workgroup::thread_rank(); + } + __CG_QUALIFIER__ bool is_valid() const { + return internal::workgroup::is_valid(); + } + __CG_QUALIFIER__ void sync() const { + internal::workgroup::sync(); + } +}; + +/** \brief User exposed API interface to construct workgroup cooperative + * group type object - `thread_block` + * + * \details User is not allowed to directly construct an object of type + * `thread_block`. Instead, he should construct it through this API + * function + */ +__CG_QUALIFIER__ thread_block +this_thread_block() { + return thread_block(internal::workgroup::size()); +} + /** * Implemenation of all publicly exposed base class APIs */ @@ -177,6 +226,9 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { case internal::cg_grid: { return (static_cast(this)->thread_rank()); } + case internal::cg_workgroup: { + return (static_cast(this)->thread_rank()); + } default: { return 0; //TODO(mahesha) } @@ -191,6 +243,9 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const { case internal::cg_grid: { return (static_cast(this)->is_valid()); } + case internal::cg_workgroup: { + return (static_cast(this)->is_valid()); + } default: { return false; } @@ -207,6 +262,10 @@ __CG_QUALIFIER__ void thread_group::sync() const { static_cast(this)->sync(); break; } + case internal::cg_workgroup: { + static_cast(this)->sync(); + break; + } } } diff --git a/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h b/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h index 9738448d94..4e10c0da72 100644 --- a/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h +++ b/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h @@ -60,7 +60,8 @@ namespace internal { typedef enum { cg_invalid, cg_multi_grid, - cg_grid + cg_grid, + cg_workgroup } group_type; /** @@ -136,6 +137,43 @@ __CG_STATIC_QUALIFIER__ void sync() { } // namespace grid +/** + * Functionalities related to `workgroup` (thread_block in CUDA terminology) + * cooperative group type + */ +namespace workgroup { + +__CG_STATIC_QUALIFIER__ dim3 group_index() { + return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y, + (uint32_t)hipBlockIdx_z)); +} + +__CG_STATIC_QUALIFIER__ dim3 thread_index() { + return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y, + (uint32_t)hipThreadIdx_z)); +} + +__CG_STATIC_QUALIFIER__ uint32_t size() { + return((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); +} + +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + + (hipThreadIdx_y * hipBlockDim_x) + + (hipThreadIdx_x))); +} + +__CG_STATIC_QUALIFIER__ bool is_valid() { + //TODO(mahesha) any functionality need to be added here? I believe not + return true; +} + +__CG_STATIC_QUALIFIER__ void sync() { + __syncthreads(); +} + +} // namespace workgroup + } // namespace internal } // namespace cooperative_groups diff --git a/hipamd/tests/src/cg/hipCGThreadBlockType.cpp b/hipamd/tests/src/cg/hipCGThreadBlockType.cpp new file mode 100644 index 0000000000..ab9492c609 --- /dev/null +++ b/hipamd/tests/src/cg/hipCGThreadBlockType.cpp @@ -0,0 +1,164 @@ +/* +Copyright (c) 2020 - present 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include "test_common.h" +#include "hip/hip_cooperative_groups.h" + +#define HIP_ASSERT(lhs, rhs) assert(lhs == rhs) + +using namespace cooperative_groups; + +static __global__ +void kernel_cg_thread_block_type(dim3 *groupIndexD, + dim3 *thdIndexD, + int *sizeD, + int *thdRankD, + int *isValidD, + int *syncValD) +{ + thread_block tb = this_thread_block(); + + // Consider the workgroup id (0, 1, 1) and thread id (0, 1, 2) for validation + // of the test + int isBlockIdx011 = + (hipBlockIdx_x == 0 && hipBlockIdx_y == 1 && hipBlockIdx_z == 1); + int isThreadIdx012 = + (hipThreadIdx_x == 0 && hipThreadIdx_y == 1 && hipThreadIdx_z == 2); + + if (isBlockIdx011 && isThreadIdx012) { + *groupIndexD = tb.group_index(); + *thdIndexD = tb.thread_index(); + *sizeD = tb.size(); + *thdRankD = tb.thread_rank(); + *isValidD = tb.is_valid(); + } + + // Consider local thread id (0, 0, 0) and (0, 0, 1) for validation of sync() + // api + __shared__ int sVar[2]; + int isThreadIdx000 = + (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 0); + int isThreadIdx001 = + (hipThreadIdx_x == 0 && hipThreadIdx_y == 0 && hipThreadIdx_z == 1); + + if (isThreadIdx000) + sVar[0] = 10; + if (isThreadIdx001) + sVar[1] = 20; + tb.sync(); + if (isBlockIdx011 && isThreadIdx012) + *syncValD = sVar[0] + sVar[1]; +} + +static void test_cg_thread_block_type() +{ + dim3 *groupIndexD, *groupIndexH; + dim3 *thdIndexD, *thdIndexH; + int *sizeD, *sizeH; + int *thdRankD, *thdRankH; + int *isValidD, *isValidH; + int *syncValD, *syncValH; + + // Allocate device memory + HIP_ASSERT(hipMalloc((void**)&groupIndexD, sizeof(dim3)), hipSuccess); + HIP_ASSERT(hipMalloc((void**)&thdIndexD, sizeof(dim3)), hipSuccess); + HIP_ASSERT(hipMalloc((void**)&sizeD, sizeof(int)), hipSuccess); + HIP_ASSERT(hipMalloc((void**)&thdRankD, sizeof(int)), hipSuccess); + HIP_ASSERT(hipMalloc((void**)&isValidD, sizeof(int)), hipSuccess); + HIP_ASSERT(hipMalloc((void**)&syncValD, sizeof(int)), hipSuccess); + + // Allocate host memory + HIP_ASSERT(hipHostMalloc((void**)&groupIndexH, sizeof(dim3)), hipSuccess); + HIP_ASSERT(hipHostMalloc((void**)&thdIndexH, sizeof(dim3)), hipSuccess); + HIP_ASSERT(hipHostMalloc((void**)&sizeH, sizeof(int)), hipSuccess); + HIP_ASSERT(hipHostMalloc((void**)&thdRankH, sizeof(int)), hipSuccess); + HIP_ASSERT(hipHostMalloc((void**)&isValidH, sizeof(int)), hipSuccess); + HIP_ASSERT(hipHostMalloc((void**)&syncValH, sizeof(int)), hipSuccess); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_thread_block_type, + dim3(2, 2, 2), + dim3(4, 4, 4), + 0, + 0, + groupIndexD, + thdIndexD, + sizeD, + thdRankD, + isValidD, + syncValD); + + // Copy result from device to host + HIP_ASSERT(hipMemcpy(groupIndexH, groupIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess); + HIP_ASSERT(hipMemcpy(thdIndexH, thdIndexD, sizeof(dim3), hipMemcpyDeviceToHost), hipSuccess); + HIP_ASSERT(hipMemcpy(sizeH, sizeD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); + HIP_ASSERT(hipMemcpy(thdRankH, thdRankD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); + HIP_ASSERT(hipMemcpy(isValidH, isValidD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); + HIP_ASSERT(hipMemcpy(syncValH, syncValD, sizeof(int), hipMemcpyDeviceToHost), hipSuccess); + + // Validate result + // Group index should be (0, 1, 1) + HIP_ASSERT(groupIndexH->x, 0); + HIP_ASSERT(groupIndexH->y, 1); + HIP_ASSERT(groupIndexH->z, 1); + // Thread index should be (0, 1, 2) + HIP_ASSERT(thdIndexH->x, 0); + HIP_ASSERT(thdIndexH->y, 1); + HIP_ASSERT(thdIndexH->z, 2); + // Workgroup size should be 64 + HIP_ASSERT(*sizeH, 64); + // Thread rank should be 36 + HIP_ASSERT(*thdRankH, 36); + // Call to is_valid() should return true + HIP_ASSERT(*isValidH, 1); + // syncVal should be equal to 30 + HIP_ASSERT(*syncValH, 30); + + // Free device memory + HIP_ASSERT(hipFree(groupIndexD), hipSuccess); + HIP_ASSERT(hipFree(thdIndexD), hipSuccess); + HIP_ASSERT(hipFree(sizeD), hipSuccess); + HIP_ASSERT(hipFree(thdRankD), hipSuccess); + HIP_ASSERT(hipFree(isValidD), hipSuccess); + HIP_ASSERT(hipFree(syncValD), hipSuccess); + + //Free host memory + HIP_ASSERT(hipHostFree(groupIndexH), hipSuccess); + HIP_ASSERT(hipHostFree(thdIndexH), hipSuccess); + HIP_ASSERT(hipHostFree(sizeH), hipSuccess); + HIP_ASSERT(hipHostFree(thdRankH), hipSuccess); + HIP_ASSERT(hipHostFree(isValidH), hipSuccess); + HIP_ASSERT(hipHostFree(syncValH), hipSuccess); +} + +int main() +{ + test_cg_thread_block_type(); + passed(); +} From ee4688d37e4796b30011c12716ff53039bffaf95 Mon Sep 17 00:00:00 2001 From: Dittakavi Satyanvesh Date: Thu, 21 May 2020 06:07:57 -0400 Subject: [PATCH 7/9] hipIpcCloseMemHandle checks the status of IpcDetach Change-Id: Ifbe8e5bbda610a1007f881627d0da1c874d03682 --- hipamd/rocclr/hip_memory.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index 3aad7f57ba..be35c40ad8 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/hipamd/rocclr/hip_memory.cpp @@ -2023,7 +2023,9 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) { amd::MemObjMap::RemoveMemObj(amd_mem_obj); /* detach the memory */ - device->IpcDetach(*amd_mem_obj); + if (!device->IpcDetach(*amd_mem_obj)){ + HIP_RETURN(hipErrorInvalidHandle); + } HIP_RETURN(hipSuccess); } From 6bc01b31d68a04f6a4b4d12ccfbf5e61d032b5a6 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Fri, 22 May 2020 18:06:37 -0400 Subject: [PATCH 8/9] SWDEV - 237467 - Return proper hip error codes incase of ROCclr IPC API failures. Change-Id: I2cc8da543f70bb3d8b82520fa9b2f509d20ce3c0 --- hipamd/rocclr/hip_memory.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index be35c40ad8..b7b8d77592 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/hipamd/rocclr/hip_memory.cpp @@ -1962,7 +1962,10 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) { /* Create an handle for IPC. Store the memory size inside the handle */ ihandle = reinterpret_cast(handle); - dev_mem_obj->IpcCreate(offset, &(ihandle->psize), &(ihandle->ipc_handle)); + if(!dev_mem_obj->IpcCreate(offset, &(ihandle->psize), &(ihandle->ipc_handle))) { + DevLogPrintfError("IPC memory creation failed for memory: 0x%x", dev_ptr); + HIP_RETURN(hipErrorInvalidValue); + } HIP_RETURN(hipSuccess); } From ed11059230f9846fdd83ad6cf777be58f2f4ab8c Mon Sep 17 00:00:00 2001 From: Evgeny Date: Fri, 22 May 2020 23:50:59 -0500 Subject: [PATCH 9/9] hip_prof: fixing printing pointer args Change-Id: I93969723650f7c29d5c00a3809d3701c6a3dca44 --- hipamd/rocclr/hip_prof_gen.py | 22 +++++++++++++++++----- 1 file changed, 17 insertions(+), 5 deletions(-) diff --git a/hipamd/rocclr/hip_prof_gen.py b/hipamd/rocclr/hip_prof_gen.py index c508888140..c59d1cca35 100755 --- a/hipamd/rocclr/hip_prof_gen.py +++ b/hipamd/rocclr/hip_prof_gen.py @@ -452,7 +452,12 @@ def generate_prof_header(f, api_map, opts_map): arg_type = arg_tuple[0] fld_name = arg_tuple[1] arg_name = opts_list[ind] - f.write(' cb_data.args.' + name + '.' + fld_name + ' = (' + arg_type + ')' + arg_name + '; \\\n') + if arg_type == "char*": + f.write(' cb_data.args.' + name + '.' + fld_name + ' = strdup("***"); \\\n') + elif arg_type == "const char*": + f.write(' cb_data.args.' + name + '.' + fld_name + ' = (' + arg_name + ') ? strdup(' + arg_name + ') : NULL; \\\n') + else: + f.write(' cb_data.args.' + name + '.' + fld_name + ' = (' + arg_type + ')' + arg_name + '; \\\n') f.write('};\n') f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n') @@ -467,13 +472,20 @@ def generate_prof_header(f, api_map, opts_map): f.write(' switch (id) {\n') for name, args in api_map.items(): f.write(' case HIP_API_ID_' + name + ':\n') - f.write(' oss << "' + name + '("') + f.write(' oss << "' + name + '(";\n') for ind in range(0, len(args)): arg_tuple = args[ind] + arg_type = arg_tuple[0] arg_name = arg_tuple[1] - if ind != 0: f.write(' << ","') - f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name) - f.write('\n << ")";\n') + var_name = 'data->args.' + name + '.' + arg_name + delim = '' if ind == 0 else ', '; + oss_stream = 'oss << "' + delim + arg_name + '=' + line_shift = ' ' + f.write(line_shift) + if re.search(r'\*', arg_type): + f.write('if (' + var_name + ' == NULL) ' + oss_stream + 'NULL";\n' + line_shift + 'else ') + f.write(oss_stream + '" << ' + var_name + ';\n') + f.write(' oss << ")";\n') f.write(' break;\n') f.write(' default: oss << "unknown";\n') f.write(' };\n')