Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Merge 6.2 hotfixes into 6.3 (#607) #620

Merged
merged 1 commit into from
Oct 11, 2024
Merged

Merge 6.2 hotfixes into 6.3 (#607) #620

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 6 additions & 1 deletion CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,12 @@ Documentation for rocPRIM is available at

* `rocprim::thread_load` and `rocprim::thread_store`, use dereference instead. Not all of those functions are available on every device architecture, and their usage can hurt performance, because inline assembly inhibits optimizations.

## Unreleased rocPRIM-3.2.0 for ROCm 6.2.0
## rocPRIM-3.2.1 for ROCm 6.2.1

### Optimizations
* Improved performance of block_reduce_warp_reduce when warp size == block size.

## rocPRIM-3.2.0 for ROCm 6.2.0

### Additions

Expand Down
15 changes: 13 additions & 2 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,6 +74,11 @@ set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CXX_EXTENSIONS OFF)

if(DEFINED BUILD_SHARED_LIBS)
set(PKG_BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS})
else()
set(PKG_BUILD_SHARED_LIBS ON)
endif()
set(BUILD_SHARED_LIBS OFF) # don't build client dependencies as shared
if(NOT USE_HIP_CPU)
# Get dependencies (required here to get rocm-cmake)
Expand Down Expand Up @@ -172,16 +177,22 @@ if(BUILD_DOCS AND NOT ONLY_INSTALL)
add_subdirectory(docs)
endif()

# set BUILD_SHARED_LIBS for packaging
set(BUILD_SHARED_LIBS ${PKG_BUILD_SHARED_LIBS})
# Package
if (ROCPRIM_PROJECT_IS_TOP_LEVEL)
set(BUILD_SHARED_LIBS ON) # Build as though shared library for naming
# add dependency on HIP runtime
set(HIP_RUNTIME_MINIMUM 4.5.0)
if(BUILD_ADDRESS_SANITIZER)
set(DEPENDS_HIP_RUNTIME "hip-runtime-amd-asan" )
else()
set(DEPENDS_HIP_RUNTIME "hip-runtime-amd" )
endif()

rocm_package_add_dependencies(DEPENDS "${DEPENDS_HIP_RUNTIME} >= 4.5.0")
rocm_package_add_dependencies(SHARED_DEPENDS "${DEPENDS_HIP_RUNTIME} >= ${HIP_RUNTIME_MINIMUM}")
rocm_package_add_deb_dependencies(STATIC_DEPENDS "hip-static-dev >= ${HIP_RUNTIME_MINIMUM}")
rocm_package_add_rpm_dependencies(STATIC_DEPENDS "hip-static-devel >= ${HIP_RUNTIME_MINIMUM}")

set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.txt")
set(CPACK_RPM_PACKAGE_LICENSE "MIT")

Expand Down
10 changes: 5 additions & 5 deletions docs/sphinx/requirements.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ docutils==0.21.2
# myst-parser
# pydata-sphinx-theme
# sphinx
fastjsonschema==2.19.1
fastjsonschema==2.20.0
# via rocm-docs-core
gitdb==4.0.11
# via gitpython
Expand All @@ -62,13 +62,13 @@ mdurl==0.1.2
# via markdown-it-py
myst-parser==3.0.1
# via rocm-docs-core
packaging==24.0
packaging==24.1
# via
# pydata-sphinx-theme
# sphinx
pycparser==2.22
# via cffi
pydata-sphinx-theme==0.15.3
pydata-sphinx-theme==0.15.4
# via
# rocm-docs-core
# sphinx-book-theme
Expand Down Expand Up @@ -111,7 +111,7 @@ sphinx==7.3.7
# sphinx-design
# sphinx-external-toc
# sphinx-notfound-page
sphinx-book-theme==1.1.2
sphinx-book-theme==1.1.3
# via rocm-docs-core
sphinx-copybutton==0.5.2
# via rocm-docs-core
Expand All @@ -135,7 +135,7 @@ sphinxcontrib-serializinghtml==1.1.10
# via sphinx
tomli==2.0.1
# via sphinx
typing-extensions==4.12.0
typing-extensions==4.12.2
# via
# pydata-sphinx-theme
# pygithub
Expand Down
66 changes: 37 additions & 29 deletions rocprim/include/rocprim/block/detail/block_reduce_warp_reduce.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,21 +180,25 @@ class block_reduce_warp_reduce
input, output, num_valid, reduce_op
);

// i-th warp will have its partial stored in storage_.warp_partials[i-1]
if(lane_id == 0)
// Final reduction across warps is only required if there is more than 1 warp
if ROCPRIM_IF_CONSTEXPR (warps_no_ > 1)
{
storage_.warp_partials[warp_id] = output;
}
::rocprim::syncthreads();

if(flat_tid < warps_no_)
{
// Use warp partial to calculate the final reduce results for every thread
auto warp_partial = storage_.warp_partials[lane_id];

warp_reduce<!warps_no_is_pow_of_two_, warp_reduce_output_type>(
warp_partial, output, warps_no_, reduce_op
);
// i-th warp will have its partial stored in storage_.warp_partials[i-1]
if(lane_id == 0)
{
storage_.warp_partials[warp_id] = output;
}
::rocprim::syncthreads();

if(flat_tid < warps_no_)
{
// Use warp partial to calculate the final reduce results for every thread
auto warp_partial = storage_.warp_partials[lane_id];

warp_reduce<!warps_no_is_pow_of_two_, warp_reduce_output_type>(
warp_partial, output, warps_no_, reduce_op
);
}
}
}

Expand Down Expand Up @@ -246,22 +250,26 @@ class block_reduce_warp_reduce
input, output, num_valid, reduce_op
);

// i-th warp will have its partial stored in storage_.warp_partials[i-1]
if(lane_id == 0)
// Final reduction across warps is only required if there is more than 1 warp
if ROCPRIM_IF_CONSTEXPR (warps_no_ > 1)
{
storage_.warp_partials[warp_id] = output;
}
::rocprim::syncthreads();

if(flat_tid < warps_no_)
{
// Use warp partial to calculate the final reduce results for every thread
auto warp_partial = storage_.warp_partials[lane_id];

unsigned int valid_warps_no = (valid_items + warp_size_ - 1) / warp_size_;
warp_reduce_output_type().reduce(
warp_partial, output, valid_warps_no, reduce_op
);
// i-th warp will have its partial stored in storage_.warp_partials[i-1]
if(lane_id == 0)
{
storage_.warp_partials[warp_id] = output;
}
::rocprim::syncthreads();

if(flat_tid < warps_no_)
{
// Use warp partial to calculate the final reduce results for every thread
auto warp_partial = storage_.warp_partials[lane_id];

unsigned int valid_warps_no = (valid_items + warp_size_ - 1) / warp_size_;
warp_reduce_output_type().reduce(
warp_partial, output, valid_warps_no, reduce_op
);
}
}
}
};
Expand Down