Skip to content

Conversation

@umfranzw
Copy link
Collaborator

We recently added a "fallback" implementation for thrust::copy_if that is invoked when copying a custom type that's too large to fit in shared memory.

This change extends the fallback slightly so that it can be used with an overload of copy_if that accepts a stencil buffer (to copy by key). It also adds a unit test to cover this case.

It also fixes a small bug in the fallback implementation that could cause the scan accumulator type to overflow when the results are compacted.

Copy link
Member

@Naraenda Naraenda left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the changes. Tests and refactor look good. I have one small question regarding the performance impact of changing the type of the flags and another remark regarding a possible optimisation.

We recently added a "fallback" implementation for thrust::copy_if that
is invoked when copying a custom type that's too large to fit in
shared memory.

This change extends the fallback slightly so that it can be used with
an overload of copy_if that accepts a stencil buffer (to copy by
key). It also adds a unit test to cover this case.

It also fixes a small bug in the fallback implementation that could
cause the scan accumulator type to overflow when the results are
compacted.
@umfranzw umfranzw merged commit c806fb5 into ROCm:release-staging/rocm-rel-6.4 Jan 20, 2025
13 of 20 checks passed
umfranzw added a commit that referenced this pull request Feb 26, 2025
* Enable HIP language (#493)

If the HIP language can be used, the HIP_USECXX variable sets rocThrust to use the CMake HIP language rather than CXX.

Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>

* Made TBB optional for hipstdpar tests (#507)

* added TBB in dependencies cmake

* updated changelog

* removed duplicate dependencies declaration in test cmake

* Create optional flag to run tests hipstdpar with TBB

* Add TBB optional flag to CHANGELOG

* Add rocprim location to hipstdpar

* Add default and test for DONWLOAD_ROCRAND

* Build with TBB when available or when FLAG is set

---------

Co-authored-by: NguyenNhuDi <Zee.nguyen@amd.com>

* Separate CMake BUILD_TEST and BUILD_HIPSTDPAR_TEST options (#508)

* Separate CMake BUILD_TEST and BUILD_HIPSTDPAR_TEST options

Previously, enabling BUILD_TEST would also enable hipstdpar tests
if we detected that a c++17-capable compiler was present. However,
this caused build issues on systems with a c++17 compiler but an
outdated version of libstdc++ that didn't support c++17 (RHEL 8.x).

Currently, we require a minimum cmake version of 3.10.2. There's
no real robust way of detecting the libstdc++ version that will work
that far back.

To workaround this problem for now, this change splits the BUILD_TEST
and BUILD_HIPSTDPAR_TEST cmake options so that they are independent.
This means that in order to enable hipstdpar tests, the user must
explicitly enable the BUILD_HIPSTDPAR_TEST option.

Update the readme to reflect this.

* Update README.md

Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>

---------

Co-authored-by: Di Nguyen <dinguyennhu@gmail.com>
Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>

* Updated known issue for inclusive_scan_by_key compiler bug (#513)

* updated the known issue

* Update CHANGELOG.md

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* moved known issue to 6.4 section

* updated to include that the issue has been fixed

---------

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Extend fallback coverage for copy_if (#512)

We recently added a "fallback" implementation for thrust::copy_if that
is invoked when copying a custom type that's too large to fit in
shared memory.

This change extends the fallback slightly so that it can be used with
an overload of copy_if that accepts a stencil buffer (to copy by
key). It also adds a unit test to cover this case.

It also fixes a small bug in the fallback implementation that could
cause the scan accumulator type to overflow when the results are
compacted.

---------

Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>
Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com>
Co-authored-by: NguyenNhuDi <Zee.nguyen@amd.com>
Co-authored-by: Di Nguyen <dinguyennhu@gmail.com>
Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>
Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>
ammallya pushed a commit that referenced this pull request Oct 28, 2025
* Enable HIP language (#493)

If the HIP language can be used, the HIP_USECXX variable sets rocThrust to use the CMake HIP language rather than CXX.

Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>

* Made TBB optional for hipstdpar tests (#507)

* added TBB in dependencies cmake

* updated changelog

* removed duplicate dependencies declaration in test cmake

* Create optional flag to run tests hipstdpar with TBB

* Add TBB optional flag to CHANGELOG

* Add rocprim location to hipstdpar

* Add default and test for DONWLOAD_ROCRAND

* Build with TBB when available or when FLAG is set

---------

Co-authored-by: NguyenNhuDi <Zee.nguyen@amd.com>

* Separate CMake BUILD_TEST and BUILD_HIPSTDPAR_TEST options (#508)

* Separate CMake BUILD_TEST and BUILD_HIPSTDPAR_TEST options

Previously, enabling BUILD_TEST would also enable hipstdpar tests
if we detected that a c++17-capable compiler was present. However,
this caused build issues on systems with a c++17 compiler but an
outdated version of libstdc++ that didn't support c++17 (RHEL 8.x).

Currently, we require a minimum cmake version of 3.10.2. There's
no real robust way of detecting the libstdc++ version that will work
that far back.

To workaround this problem for now, this change splits the BUILD_TEST
and BUILD_HIPSTDPAR_TEST cmake options so that they are independent.
This means that in order to enable hipstdpar tests, the user must
explicitly enable the BUILD_HIPSTDPAR_TEST option.

Update the readme to reflect this.

* Update README.md

Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>

---------

Co-authored-by: Di Nguyen <dinguyennhu@gmail.com>
Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>

* Updated known issue for inclusive_scan_by_key compiler bug (#513)

* updated the known issue

* Update CHANGELOG.md

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* moved known issue to 6.4 section

* updated to include that the issue has been fixed

---------

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Extend fallback coverage for copy_if (#512)

We recently added a "fallback" implementation for thrust::copy_if that
is invoked when copying a custom type that's too large to fit in
shared memory.

This change extends the fallback slightly so that it can be used with
an overload of copy_if that accepts a stencil buffer (to copy by
key). It also adds a unit test to cover this case.

It also fixes a small bug in the fallback implementation that could
cause the scan accumulator type to overflow when the results are
compacted.

---------

Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com>
Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com>
Co-authored-by: NguyenNhuDi <Zee.nguyen@amd.com>
Co-authored-by: Di Nguyen <dinguyennhu@gmail.com>
Co-authored-by: Leo Paoletti <164940351+lpaoletti@users.noreply.github.com>
Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

[ROCm/rocThrust commit: 0006584]
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants