Skip to content

Navigation Menu

Appearance settings

- AI CODE CREATION
  GitHub CopilotWrite better code with AI
  GitHub Copilot appDirect agents from issue to merge
  MCP Registry^NewIntegrate external tools
- DEVELOPER WORKFLOWS
  ActionsAutomate any workflow
  CodespacesInstant dev environments
  IssuesPlan and track work
  Code ReviewManage code changes
- APPLICATION SECURITY
  GitHub Advanced SecurityFind and fix vulnerabilities
  Code securitySecure your code as you build
  Secret protectionStop leaks before they start
- EXPLORE
  Why GitHub
  Documentation
  Blog
  Changelog
  Marketplace
View all features
- BY COMPANY SIZE
  Enterprises
  Small and medium teams
  Startups
  Nonprofits
- BY USE CASE
  App Modernization
  DevSecOps
  DevOps
  CI/CD
  View all use cases
- BY INDUSTRY
  Healthcare
  Financial services
  Manufacturing
  Government
  View all industries
View all solutions
- EXPLORE BY TOPIC
  AI
  Software Development
  DevOps
  Security
  View all topics
- EXPLORE BY TYPE
  Customer stories
  Events & webinars
  Ebooks & reports
  Business insights
  GitHub Skills
- SUPPORT & SERVICES
  Documentation
  Customer support
  Community forum
  Trust center
  Partners
View all resources
- COMMUNITY
  GitHub SponsorsFund open source developers
- PROGRAMS
  Security Lab
  Maintainer Community
  Accelerator
  GitHub Stars
  Archive Program
- REPOSITORIES
  Topics
  Trending
  Collections
- ENTERPRISE SOLUTIONS
  Enterprise platformAI-powered developer platform
- AVAILABLE ADD-ONS
  GitHub Advanced SecurityEnterprise-grade security features
  Copilot for BusinessEnterprise-grade AI features
  Premium SupportEnterprise-grade 24/7 support
Pricing

Search code, repositories, users, issues, pull requests...

Search

Clear

Search syntax tips

Provide feedback

We read every piece of feedback, and take your input very seriously.

Include my email address so I can be contacted

Saved searches

Use saved searches to filter your results more quickly

Name

Query

To see all available qualifiers, see our documentation.

Appearance settings

You signed in with another tab or window. Reload to refresh your session. You signed out in another tab or window. Reload to refresh your session. You switched accounts on another tab or window. Reload to refresh your session.

Dismiss alert

Uh oh!

There was an error while loading. Please reload this page.

ROCm / rocm-systems Public

Notifications You must be signed in to change notification settings
Fork 306
Star 429

Code
Issues 359
Pull requests 634
Discussions
Actions
Projects
Security and quality
Insights

Additional navigation options

Code
Issues
Pull requests
Discussions
Actions
Projects
Security and quality
Insights

AIRUNTIME-171 - cooperative groups scan#5914

Merged

g-h-c merged 100 commits into

developROCm/rocm-systems:developfrom

users/g-h-c/AIRUNTIME-171_cooperative_groups_scanROCm/rocm-systems:users/g-h-c/AIRUNTIME-171_cooperative_groups_scanCopy head branch name to clipboard

Jun 26, 2026

Conversation Commits100 (100)Checks Files changed

Merged

AIRUNTIME-171 - cooperative groups scan#5914
g-h-c merged 100 commits into
developfrom
users/g-h-c/AIRUNTIME-171_cooperative_groups_scan

Commits

Commits on Jun 26, 2026

AIRUNTIME-171 - add initial implemention of cg::inclusive_scan. Only for 4-bytes for now
g-h-c
committed
AIRUNTIME-171 - use __ockl_wfscan_add_i32() when __OPTIMIZE__ is defined
g-h-c
committed
AIRUNTIME-171 - implement inclusive_scan for any thread block tile size (not just the warp size)
g-h-c
committed
AIRUNTIME-171 - make sure the OCKL intrinsic exist before using it
g-h-c
committed
AIRUNTIME-171 - try to avoid the cost of the conditional check to see if the OCKL warp-wide intrinsic can be used, as we know the size of the tile at compile time
g-h-c
committed
Fix assertion being triggered in __hip_check_mask() for some coooperative_groups::reduce() tests
g-h-c
committed
AIRUNTIME-171 - make calculateExpected() also calculate expected results of scans depending on a parameter. A first step to allow reduce tests code to be reused for inclusive/exclusive_scan tests
g-h-c
committed
AIRUNTIME-171 - add macro GENERATE_SCAN_FUNC() to generate the overloads that call the __ockl_wfscan_* intrinsics
g-h-c
committed
AIRUNTIME-171 - refactor reduce tests so they can be reused for inclusive/exclusive scans
g-h-c
committed
ROCM-1254 - remove Unit_Thread_Block_Tile_Reduce_Non_Participating_Threads; if a thread of a cooperative group tile block does not participate, it is undefined behaviour
g-h-c
committed
Remove unused partialSum kernel (dead code with UB pattern)

Copilot
and
g-h-c
committed
AIRUNTIME-171 - make sure the expected result in a reduce is the result associated with the last lane. Make sure both reduces and scans have the final reduce result in the last lane and not the fir…
g-h-c
committed
AIRUNTIME-171 - fix that the reduction tree implementation in calculateExpected(), being in the CPU and thusly serial, was using the values in utput[i] that were already overwritten at the beginnin…
g-h-c
committed
AIRUNTIME-171 - fix that floating point aggregations were not expecting the result in the right lane
g-h-c
committed
AIRUNTIME-171 - iterate the 'modulo' variable until nextPowerOf2(lastLane + 1); lastLane + 1 is not enough for some masks
g-h-c
committed
AIRUNTIME-171 - fix we were calling compareFloatingPoint() even for lanes that were not mentioned in the mask
g-h-c
committed
AIRUNTIME-171 - add support for cooperative_groups::less in cg::inclusive_scan
g-h-c
committed
AIRUNTIME-171 - add Unit_Thread_Block_Tile_Scan_Random_arithmetic
g-h-c
committed
AIRUNTIME-171 - extend cg::inclusive_scan() to support all types, not just int
g-h-c
committed
AIRUNTIME-171 - simplify code that invokes ockl scan intrinsics. Do not take into account HIP_ENABLE_EXTRA_WARP_SYNC_TYPES for scan
g-h-c
committed
AIRUNTIME-171 - simplify mask generation logic
g-h-c
committed
AIRUNTIME-171 - fix up RTC tests compilation
g-h-c
committed
AIRUNTIME-171 - add cg::exclusive_scan and some of the associated tests
g-h-c
committed
AIRUNTIME-171 - refactor duplicated code that calculate the group mask
g-h-c
committed
AIRUNTIME-171 - fix up previous commits
g-h-c
committed
AIRUNTIME-171 - fix some scope Catch2 INFO() macro would have no effect because they were in the same scope as REQUIRE()
g-h-c
committed
AIRUNTIME-171 - fix up calculation of the expected value for cg::bit_xor
g-h-c
committed
AIRUNTIME-171 - implement cg::exclusive_scan() via a backward permute at the end. WIP. needs refactoring to avoid duplicated code
g-h-c
committed
AIRUNTIME-171 - add more cg::exclusive_scan() tests
g-h-c
committed
AIRUNTIME-171 - fix up one compareFloatingPoint() missing the Op template parameter
g-h-c
committed
AIRUNTIME-171 - add bPermute() implementing backward permutes that can take arguments of arbitrary size, as this code is repeated throughout cg::reduce(), scans and reduce_*_sync()
g-h-c
committed
AIRUNTIME-171 - start using bPermute. Remove function isPrimitiveType()
g-h-c
committed
AIRUNTIME-171 - use ockl functions for __half cooperative_group scans too
g-h-c
committed
AIRUNTIME-171 - fix up add missing tests in cooperativeGrps.yaml. Fix compilation error
g-h-c
committed
AIRUNTIME-171 - add Unit_Thread_Block_Tile_Scan_Trivially_Copyable_Parameters
g-h-c
committed
AIRUNTIME-171 - add inclusive/exclusive_scan benchmarks
g-h-c
committed
AIRUNTIME-171 - add more information when the output of the reduce benchmarks do not match the expected
g-h-c
committed
AIRUNTIME-171 - remove ExcludeFirst template parameter as all threads with tiled groups with a compile time size need to be active
g-h-c
committed
AIRUNTIME-171 - remove mask generation duplicated code in warpSync.cc
g-h-c
committed
AIRUNTIME-171 - add Unit_Thread_Block_Tile_Scan_All_Parameter_Sizes
g-h-c
committed
AIRUNTIME-171 - remove outdated comments
g-h-c
committed
AIRUNTIME-171 - fix compilation error: std::memset() is not defined in __device__ code
g-h-c
committed
AIRUNTIME-171 - fix duplicated variable
g-h-c
committed
AIRUNTIME-171 - fix integer constant not having the right suffix for wavefronts of 64 lanes
g-h-c
committed
AIRUNTIME-171 - fix use of undeclared variable
g-h-c
committed
AIRUNTIME-171 - refactor rtc cooperative groups tests so they can be reused to test inclusive/exclusive scan
g-h-c
committed
AIRUNTIME-171 - add Unit_Rtc_CoopScan
g-h-c
committed
AIRUNTIME-171 - use 1ull instead of 1ul when dealing with 64-bit masks
g-h-c
committed
AIRUNTIME-171 - 'unsigned long warpMask' should have been 'unsigned long long warpMask'
g-h-c
committed
AIRUNTIME-171 - fix another use of 1ul when it should have 1ull. Also commas where used when ; was the intention
g-h-c
committed
AIRUNTIME-171 - fix Unit_Thread_Block_Coalesced_Scan_boolean using arithmetic operators and not boolean
g-h-c
committed
AIRUNTIME-171 - fix GENERATE_SCAN_FUNC() for fp16 when not in cooperative_groups::impl
g-h-c
committed
AIRUNTIME-171 - fix potential deadlock in applyFunctor() and applyScanFunctor() (as flagged by Github Copilot in applyScanFunctor())
g-h-c
committed
AIRUNTIME-171 - add missing semicolon

g-h-c
and
Copilot
committed
AIRUNTIM-171 - Fix another use of 1ul and not 1ull on mask calculations

g-h-c
and
Copilot
committed
AIRUNTIME-171 - fix Unit_Rtc_CoopScan failing because calculateExpected() was not returning the expected values in 'output' for all the lanes, only the last
g-h-c
committed
AIRUNTIME-171 - fix OCKL intrinsics not being called for scan
g-h-c
committed
AIRUNTIME-171 - fix OCKL boolean scan intrinsics not being called
g-h-c
committed
Add hip_scan.h to hiprtc/CMakeLists.txt
g-h-c
committed
AIRUNTIME-171 - fix that in cg::exclusive_scan the identity needs to be returned for first active lane. Fix some tests not indicating the operator after a failure
g-h-c
committed
AIRUNTIME-171 - implement __hip_internal::numeric_limits<T> using constexpr methods, to avoid having to specify __device__ as it is in host_defines.h. Only invoke GENERATE_SCAN_FUNC() macro when __…
g-h-c
committed
AIRUNTIME-171 - rename __hip_internal::numeric_limits<T>::max() to NumericLimits<T>::maximum() to avoid clash with Win32 max() macro
g-h-c
committed
AIRUNTIME-171 - fix invocation of opToString should have been: opToString<T, Op<T>>()
g-h-c
committed
AIRUNTIME-171 - fix struct NumericLimits<float>::lowest() should return the minimum finite value, and not -inf
g-h-c
committed
AIRUNTIME-171 - fix rocsparse 'error: constexpr function never produces a constant expression (-Winvalid-constexpr)'
g-h-c
committed
AIRUNTIME-171 - rename opToString() in rtc_reduce.cc to reduceOpToString(), to not confuse it with the one in warp_common.hh
g-h-c
committed
AIRUNTIME-171 - fix that Unit_Rtc_CoopScan for cg::plus<float> used equality comparison, when some tolerance should be used instead in the comparison
g-h-c
committed
AIRUNTIME-171 - fix that NumericLimits<__half>::maximum() and lowest() need to be constexpr to be used in __device__ code without problem. Use __builtin_bit_cast() instead to try to allow the compi…
g-h-c
committed
AIRUNTIME-171 - fix 'expected' variable assigned but never used
g-h-c
committed
AIRUNTIME-171 - remove halfWaveSize which is actually shadowed by another in initializeMasks()
g-h-c
committed
AIRUNTIME-171 - fix undefined behaviour that could happen if __builtin_clzll() is called with an empty mask
g-h-c
committed
AIRUNTIME-171 - try to implement NumericLimits<__half>::maximum() and lowest() as constexpr using __half_raw
g-h-c
committed
AIRUNTIME-171 - avoid 'error: change of the active member of a union from '__half_raw::<unnamed union>::data' to '__half_raw::<unnamed union>::x' is not a constant expression before C++20
g-h-c
committed
AIRUNTIME-171 - fix expected value for __half in exclusive_scans tests
g-h-c
committed
AIRUNTIME-171 - execute Unit_Thread_Block_Tile_Exclusive_Scan_Basic for 'half' and cooperative_groups::less too, to make sure the first active lane returns the identity in that case
g-h-c
committed
AIRUNTIME-171 - fix up previous commit, std::numeric_limits only needs to be used for cooperative_groups::less
g-h-c
committed
AIRUNTIME-171 - fix value for NumericLimits<float>::lowest()
g-h-c
committed
AIRUNTIME-171 - fix Unit_Rtc_CoopReduce being too verbose
g-h-c
committed
AIRUNTIME-171 - make sure the operands and the return value are aligned at the right alignment in cooperative_groups operations
g-h-c
committed
AIRUNTIME-171 - prevent memcpy() in cooperative_groups::scan() to read past variables of type Val
g-h-c
committed
AIRUNTIME-171 - fix memcpy() reading past Val also for reduce operations
g-h-c
committed
AIRUNTIME-171 - make cooperative_groups::exclusive_scan() results for the first active lane to mimic what ockl_wfscan does
g-h-c
committed
Fix expected values for cg::exlusive_scan() for the first active lane
g-h-c
committed
AIRUNTIME-171 - cosmetic changes
g-h-c
committed
AIRUNTIME-171 - rename GENERATE_SCAN_FUNC as HIP_IMPL_GENERATE_SCAN_FUNC to emphasize that is a private AMD implementation detail
g-h-c
committed
AIRUNTIME-171 - fix typo in cooperativeGrps.yaml
g-h-c
committed
AIRUNTIME-171 - fix Max::operator() contained an unnecessary extra loop
g-h-c
committed
AIRUNTIME-171 - fix whitespace errors
g-h-c
committed
AIRUNTIME-171 - add missing nvidia_hip_cooperative_groups_scan.h
g-h-c
committed
AIRUNTIME-171 - add cg::inclusive_scan/exclusive_scan() overloads that do not take the operator
g-h-c
committed
AIRUNTIME-171 - fix Unit_Thread_Block_Tile_Inclusive_Scan_Basic failures after introducing AggregationType::InclusiveScanDefault and ExclusiveScanDefault
g-h-c
committed
AIRUNTIME-171 - fix that cooperative_groups/scan.h would not compile unless cooperative_groups/reduce.h was included earlier
g-h-c
committed
AIRUNTIME-171 - add test Unit_Thread_Block_Coalesced_Scan_Partition
g-h-c
committed
AIRUNTIME-171 - fix that scan should use coalesced_info.member_mask and not __activemask()
g-h-c
committed
AIRUNTIME-171 - add a test for partitioned thread_block_tiles too
g-h-c
committed
AIRUNTIME-171 - fix order of parameters when applying scan operations. Add test Unit_Thread_Block_Tile_Scan_Custom_Op which tests this behaviour using a non-commutative operator
g-h-c
committed
AIRUNTIME-171 - return {} as identity for custom types as the result of the first active lane in exclusive_scan, to be more similar to CUDA
g-h-c
committed
AIRUNTIME-171 - change cooperative_group::impl::buildMask() to use internal::workgroup::thread_rank() as cooperative_groups operations were generating a warp mask that was dependent on threadIdx.x …
g-h-c
committed
SWDEV-515087 - replace __lane_id() with a manual calculation; otherwise the 'Build clr + hip-tests (NVIDIA)' CI run would fail to compile
g-h-c
committed
ROCM-26483 - make sure cooperative_groups::reduce() will now work with 2D or 3D tiles
g-h-c
committed

Footer

© 2026 GitHub, Inc.

Footer navigation

Terms
Privacy
Security
Status
Community
Docs
Contact

You can’t perform that action at this time.