Code

Develop stream 2024-06-26#575

Code

Merged

stanleytsang-amd merged 174 commits intoROCm:developROCm/rocPRIM:developfrom StreamHPC:develop_stream_20240626StreamHPC/rocPRIM:develop_stream_20240626Copy head branch name to clipboardAug 3, 2024

+22,760-5,015

Commits

Commits on Aug 2, 2024

• feat(device_transform): add tuning benchmarks and config generation for device transform

  

  Naraenda

  authored and

  Beanavil

  committed
  3c0d51e

• perf(device_transform): tuned device transform algorithm for better performance

  

  Naraenda

  authored and

  Beanavil

  committed
  58a8b02

• docs(changelog.md): add 'device_transform' improvements to changelog

  

  Naraenda

  authored and

  Beanavil

  committed
  eb18249

• feat(ConfigAutotuneSettings.cmake): allow benchmark_device_transform to tune for more block sizes

  

  Naraenda

  authored and

  Beanavil

  committed
  64e1fe2

• fix(benchmark_device_transform.cpp): fix unused type warning when compiling tuning benchmarks

  

  Naraenda

  authored and

  Beanavil

  committed
  0f8c746

• perf(device_transform.hpp): updated configs for device transform which uses a wider range of block sizes

  

  Naraenda

  authored and

  Beanavil

  committed
  62ef870

• fix(transform_config_template): added missing '::type' in general case of transform config

  

  Naraenda

  authored and

  Beanavil

  committed
  0aa581d

• refactor(benchmark_device_transform.cpp): remove duplicated code with 'benchmark_device_transform.parallel.hpp'

  

  Naraenda

  authored and

  Beanavil

  committed
  a6f7784

• docs(changelog.md): removed 'slightly' in device transform performance improvements

  

  Naraenda

  authored and

  Beanavil

  committed
  d62a8af

• fix(benchmark_device_transform): fix various build errors and warnings

  

  Naraenda

  authored and

  Beanavil

  committed
  fda70fd

• test(test_device_batch_memcpy.cpp): add simple batch copy test

  Show description for 1cf9a23
  

  Naraenda

  authored and

  Beanavil

  committed
  1cf9a23

• fix(device_batch_memcpy.hpp): use dereference instead of 'rocprim::thread_load/store'

  Show description for e7c11a7
  

  Naraenda

  authored and

  Beanavil

  committed
  e7c11a7

• revert test(test_device_batch_memcpy.cpp): add simple batch copy test

  Show description for 94f1b7a
  

  Naraenda

  authored and

  Beanavil

  committed
  94f1b7a

• test(benchmark_device_adjacent_difference.cpp): increased the default size of input so that in place uint8 benchmarks don't fit in L3 cache on select architectures

  

  Naraenda

  authored and

  Beanavil

  committed
  9222c1a

• docs(changelog.md): update changelog with benchmark changes

  

  Naraenda

  authored and

  Beanavil

  committed
  52592ca

• Added overload for match_any

  

  jblok27

  authored and

  Beanavil

  committed
  484d7c5

• Replaced section with match_any() call

  

  jblok27

  authored and

  Beanavil

  committed
  08d1641

• Fixed copyright date

  

  jblok27

  authored and

  Beanavil

  committed
  d6ce082

• Fixed formatting

  

  jblok27

  authored and

  Beanavil

  committed
  45647a6

• change match_any to runtime dispatch

  

  jblok27

  authored and

  Beanavil

  committed
  74be978

• docs(intrinsics/warp): name the correct label_bits in match_any documentation

  

  Gergely Meszaros

  authored and

  Beanavil

  committed
  c61a05f

• unified wavefront definition

  

  nolmoonen

  authored and

  Beanavil

  committed
  8a47190

• build: Remove force-inline workaround on windows

  Show description for a6ea925
  

  Gergely Meszaros

  authored and

  Beanavil

  committed
  a6ea925

• ci: enable debug builds on windows

  Show description for d2343a2
  

  Gergely Meszaros

  authored and

  Beanavil

  committed
  d2343a2

• docs: Add CHANGELOG for removing force-inline workaround

  

  Gergely Meszaros

  authored and

  Beanavil

  committed
  2f5b048

• fix clang format

  

  parbenc

  authored and

  Beanavil

  committed
  59cbb43

• fix(tests): Add saturating casts and use them for random data generation

  Show description for ab61ca9
  

  Gergely Meszaros

  authored and

  Beanavil

  committed
  ab61ca9

• fix(benchmark_device_adjacent_difference): fixe size in bytes instead of number of elements

  Beanavil

  committed
  557dd5e

• Update contributing guidelines

  Beanavil

  committed
  02fd1fe

• specify benchmark seed via command line

  

  nolmoonen

  authored and

  Beanavil

  committed
  6c4bcce

• refactor lookback sleep dispatch

  

  nolmoonen

  authored and

  Beanavil

  committed
  d7421d4

• add config to tests

  

  nolmoonen

  authored and

  Beanavil

  committed
  d9dcd9a

• add config tuning for partition

  

  nolmoonen

  authored and

  Beanavil

  committed
  691dc37

• generic tuning

  

  nolmoonen

  authored and

  Beanavil

  committed
  3e8f81f

• add tuned configurations

  

  nolmoonen

  authored and

  Beanavil

  committed
  e08b551

• Fix "warning: loop not unrolled" with CMAKE_BUILD_TYPE=MinSizeRel (-Os)

  Show description for 19ac4a6
  

  ex-rzr

  authored and

  Beanavil

  committed
  19ac4a6

• fix(device_partition): re-added workaround for the device_partition family to properly limit block size for the base configuration

  

  mfep

  authored and

  Beanavil

  committed
  93827f4

• Using .lint:clang-format

  

  mfep

  authored and

  Beanavil

  committed
  4b404e1

• refactor(intrinsics/thread.hpp): remove 'memory_fence_device' workaround for compiler bug on gfx10 and gfx11

  

  Naraenda

  authored and

  Beanavil

  committed
  36704ce

• ci(.gitlab-ciy.yml): disable debug builds in cmake-minimum due to excessive build times when targeting debug test

  

  Naraenda

  authored and

  Beanavil

  committed
  0d62eda

• fix(docs): Fixed documentation for thread subdir

  

  mfep

  authored and

  Beanavil

  committed
  fe15107

• fix(docs): Fixed documentation for the types subdir

  

  mfep

  authored and

  Beanavil

  committed
  ada912c

• fix undefined behavior in test data generation

  

  nolmoonen

  authored and

  Beanavil

  committed
  ec3f614

• Deprecate thread_load/thread_store

  

  NB4444

  authored and

  Beanavil

  committed
  a71833e

• Ignore thread_load and thread_store deprecation warnings

  

  NB4444

  authored and

  Beanavil

  committed
  3b5d460

• Deprecated raw_storage and replaced by uninitialized_array in a few locations

  

  mfep

  authored and

  Beanavil

  committed
  20d84ed

• unsigned char storage in raw_storage to prevent undefined aliasing

  

  mfep

  authored and

  Beanavil

  committed
  392f09b

• Added ROCPRIM_DONT_SUPPRESS_DEPRECATIONS

  

  mfep

  authored and

  Beanavil

  committed
  139993c

• Resolve "Improve rocPRIM test logs"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  2d07e14

• improve documentation for configuration tuning

  

  nolmoonen

  authored and

  Beanavil

  committed
  cc9f2ba

• Refactor device_scan, use is_sleep_scan_state_used and with_scan_state as in other lookback algorithms

  

  ex-rzr

  authored and

  Beanavil

  committed
  0bb05fd

• Use device of the current stream in is_sleep_scan_state_used

  

  ex-rzr

  authored and

  Beanavil

  committed
  e72d33f

• Do not build kernels with sleep in lookback state on devices that don't need it (!=gfx908)

  

  ex-rzr

  authored and

  Beanavil

  committed
  6e2681a

• Resolve "Add thread headers to rocprim.hpp and document thread-level methods"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  f41f1cf

• Resolve "Batch memcpy: disable BENCHMARK_BATCH_MEMCPY_NAIVE"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  48c7c58

• Resolve "Fix under- and overflow in minimum and maximum for input data for benchmarks"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  eea8806

• Resolve "CMake build consistency"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  a24ddf6

• Resolve "Benchmark utility for random segments generates segments of wrong size"

  

  Arsalan Anwari

  authored and

  Beanavil

  committed
  68620f4

• Adapt device segmented_reduce for large indices within a segment

  Beanavil

  committed
  67cdd25

• Add large indices test

  Beanavil

  committed
  ede8b05

• Update CHANGELOG

  Beanavil

  committed
  5735213

• reduce by key tuning

  

  nolmoonen

  authored and

  Beanavil

  committed
  eac62d4

• First commit nth element

  

  NB4444

  authored and

  Beanavil

  committed
  28a7530

• Tests nth element

  

  NB4444

  authored and

  Beanavil

  committed
  1ae80f7

• Simplified working version nth element on one block

  

  NB4444

  authored and

  Beanavil

  committed
  6878eb4

• Added output check for correctness

  

  NB4444

  authored and

  Beanavil

  committed
  9d7fd29

• nth element sizes larger then 64

  

  NB4444

  authored and

  Beanavil

  committed
  3bde5aa

• Added equality buckets to nth_element logic

  

  NB4444

  authored and

  Beanavil

  committed
  58e7a66

• Added multiple blocks for nth element

  

  NB4444

  authored and

  Beanavil

  committed
  5845cc4

• Added test to see if elements did not change

  

  NB4444

  authored and

  Beanavil

  committed
  976e6fe

• Debugging synchronization

  

  NB4444

  authored and

  Beanavil

  committed
  4d01270

• Nth element working version only for key with comperator greater and

  Show description for 3674577
  

  NB4444

  authored and

  Beanavil

  committed
  3674577

• Nth element implemented for key with tests

  

  NB4444

  authored and

  Beanavil

  committed
  978e81d

• Fixed issue for custom types in nth element and added tests

  

  NB4444

  authored and

  Beanavil

  committed
  8d26d5d

• Added input and output itterators for nth element

  

  NB4444

  authored and

  Beanavil

  committed
  f3cc3dc

• Added some benchmarks for nth element

  

  NB4444

  authored and

  Beanavil

  committed
  32356e4

• Small optimizations nth element

  

  NB4444

  authored and

  Beanavil

  committed
  e9cc712

• Debug code nth element

  

  NB4444

  authored and

  Beanavil

  committed
  fa514c0

• Made seperate kernel for block offset calculations nth element

  

  NB4444

  authored and

  Beanavil

  committed
  459e205

• Small optimizations nth element

  

  NB4444

  authored and

  Beanavil

  committed
  1146b68

• Moved all block offset calculation to other kernel nth element

  

  NB4444

  authored and

  Beanavil

  committed
  2ee4b9d

• Optimization nth element

  

  NB4444

  authored and

  Beanavil

  committed
  65ea5b9

• Make use of radix_rank instead of multiple scans

  

  NB4444

  authored and

  Beanavil

  committed
  53b1dee

• Start of adding multiple items per thread nth element

  

  NB4444

  authored and

  Beanavil

  committed
  75d15cb

• Nth element using less shared memory

  

  NB4444

  authored and

  Beanavil

  committed
  b474a4f

• Nth element small optimizations and cleanup code

  

  NB4444

  authored and

  Beanavil

  committed
  7403a43

• Fixed benchmark break nth element after rebase

  

  NB4444

  authored and

  Beanavil

  committed
  d9863c6

• nth element local oracle for buckets_store

  

  NB4444

  authored and

  Beanavil

  committed
  37ec173

• Cleanup nth element

  

  NB4444

  authored and

  Beanavil

  committed
  5ef6cc9

• Nth element update tests with random nth element

  

  NB4444

  authored and

  Beanavil

  committed
  75f58ea

• Addition of configs for nth element

  

  NB4444

  authored and

  Beanavil

  committed
  ea383b7

• Add lookbackstates to nth element

  

  NB4444

  authored and

  Beanavil

  committed
  20dc554

• Cleanup and extra comments in nth_element

  

  NB4444

  authored and

  Beanavil

  committed
  d69b533

• Removed unnecesarry test cases and choose nth_element based on seed_value

  

  NB4444

  authored and

  Beanavil

  committed
  d3c9896

• Added nth_element to changelog

  

  NB4444

  authored and

  Beanavil

  committed
  a4726a4

• Updated benchmark of nth_element based on feedback

  

  NB4444

  authored and

  Beanavil

  committed
  21c105c

• Nth_element updated tests and config based on review

  

  NB4444

  authored and

  Beanavil

  committed
  a62b86f

• Documentation updated for nth_element

  

  NB4444

  authored and

  Beanavil

  committed
  e874f16

• Cleanup code nth element

  

  NB4444

  authored and

  Beanavil

  committed
  2523f3a

• Nth element changes based on review

  

  NB4444

  authored and

  Beanavil

  committed
  e803294

• Add documentation spinx doc

  

  NB4444

  authored and

  Beanavil

  committed
  bd43291

• Changed nth element to a while loop

  

  NB4444

  authored and

  Beanavil

  committed
  50b52e9

• Nth element asserts in device code

  

  NB4444

  authored and

  Beanavil

  committed
  f8aa873

• Nth element documentation fixes

  

  NB4444

  authored and

  Beanavil

  committed
  1bacc52

• nth element docs crash fix

  

  NB4444

  authored and

  Beanavil

  committed
  d60de33

• nth element lookback state reset

  

  NB4444

  authored and

  Beanavil

  committed
  6dc847b

• Nth element changes based on review

  

  NB4444

  authored and

  Beanavil

  committed
  28afd01

• Replaced raw storage with unitialized_array in nth element

  

  NB4444

  authored and

  Beanavil

  committed
  a548d00

• Changed Nth element to be able to be used with iterators

  

  NB4444

  authored and

  Beanavil

  committed
  de43f1c

• nth element fix small mistakes

  

  NB4444

  authored and

  Beanavil

  committed
  e1bad2f

• Added config for in place nth element

  

  NB4444

  authored and

  Beanavil

  committed
  1b1e75a

• Changes based on review

  

  NB4444

  authored and

  Beanavil

  committed
  220167c

• Added c++17 tests nth_element

  

  NB4444

  authored and

  Beanavil

  committed
  56d360a

• Make use of internal merge_path also fix bug with unsigned types for size

  

  NB4444

  authored and

  Beanavil

  committed
  d78f677

• Added test for public merge_path_search

  

  NB4444

  authored and

  Beanavil

  committed
  16ca95f

• Fixed thread_load and thread_store bug with float and double

  

  NB4444

  authored and

  Beanavil

  committed
  b0a8de5

• Made review changes

  

  NB4444

  authored and

  Beanavil

  committed
  3723161

• Add bug fixes to changelog

  

  NB4444

  authored and

  Beanavil

  committed
  a027de6

• ci: remove trailing newlines in gitlab-ci.yml

  

  Snektron

  authored and

  Beanavil

  committed
  daf6a8b

• ci: compress autotune artifacts using zstd

  

  Snektron

  authored and

  Beanavil

  committed
  361c50f

• Removed oracles array from nth element

  

  NB4444

  authored and

  Beanavil

  committed
  a5b95c7

• Remove constraint of 256 for number of buckets nth element

  

  NB4444

  authored and

  Beanavil

  committed
  aba13fe

• Apply 1 suggestion(s) to 1 file(s)

  Show description for 23a3829
  

  2 peopleauthored and

  Beanavil

  committed
  23a3829

• clang-format: trick clang-format into always breaking after c-style function attributes

  

  Snektron

  authored and

  Beanavil

  committed
  30ba72b

• add ctz intrinsic

  Show description for d538a5a
  

  Snektron

  authored and

  Beanavil

  committed
  d538a5a

• lookback scan: remove HIP-CPU bits

  Show description for 9c8afb3
  

  Snektron

  authored and

  Beanavil

  committed
  9c8afb3

• lookback scan: reformat

  Show description for 15cbd89
  

  Snektron

  authored and

  Beanavil

  committed
  15cbd89

• lookback scan: add reproducibility test

  

  Snektron

  authored and

  Beanavil

  committed
  3c48cdf

• test: print floats as hexfloat in assert_bit_eq

  

  Snektron

  authored and

  Beanavil

  committed
  1dc10dc

• add warp_readfirstlane and warp_readlane intrinsics

  

  Snektron

  authored and

  Beanavil

  committed
  60502f0

• lookback scan: add deterministic implementation

  

  Snektron

  authored and

  Beanavil

  committed
  c5533d6

• scan: add deterministic overload

  

  Snektron

  authored and

  Beanavil

  committed
  6139baf

• scan_by_key: add deterministic overload

  

  Snektron

  authored and

  Beanavil

  committed
  8290563

• reduce_by_key: add deterministic overload

  

  Snektron

  authored and

  Beanavil

  committed
  ef5fd4b

• add char and short atomic load/store overloads

  Show description for 82e35ea
  

  Snektron

  authored and

  Beanavil

  committed
  82e35ea

• lookback scan: change flag to be always one byte

  Show description for 5139e95
  

  Snektron

  authored and

  Beanavil

  committed
  5139e95

• lookback scan: swap flag and prefix, allow fast scan for values up to size 7

  Show description for 4de4860
  

  Snektron

  authored and

  Beanavil

  committed
  4de4860

• nara nit f32

  

  Snektron

  authored and

  Beanavil

  committed
  db676c9

• update changelog with mention of deterministic algorithms

  

  Snektron

  authored and

  Beanavil

  committed
  258cd3e

• lookback reproducibility test: allocate temporary memory with the right scan operator

  

  Snektron

  authored and

  Beanavil

  committed
  4ea5ea6

• lookback scan: avoid caching large types

  Show description for 642c68d
  

  Snektron

  authored and

  Beanavil

  committed
  642c68d

• remove assertions in lookbacn scan, they don't compile properly in debug builds

  

  Snektron

  authored and

  Beanavil

  committed
  e1f5576

• lookback reproducibility test: use same functor for both tests

  Show description for 0ee4a1b
  

  Snektron

  authored and

  Beanavil

  committed
  0ee4a1b

• lookback scan: rotate prefix rather than block_prefix

  

  Snektron

  authored and

  Beanavil

  committed
  8cce0d3

• lookback scan: also test deterministic in normal tests

  

  Snektron

  authored and

  Beanavil

  committed
  ea287e9

• naive implementation

  

  nolmoonen

  authored and

  Beanavil

  committed
  d0f77bb

• partial sort benchmark

  

  NB4444

  authored and

  Beanavil

  committed
  c4b6566

• Made partial_sort in place and created partial_sort_copy

  

  NB4444

  authored and

  Beanavil

  committed
  e8e48be

• Add and fix documentation partial_sort

  

  NB4444

  authored and

  Beanavil

  committed
  f72156c

• Test partial_sort with iterator

  

  NB4444

  authored and

  Beanavil

  committed
  26939ae

• Add partial_sort and partial_sort_copy to the changelog

  

  NB4444

  authored and

  Beanavil

  committed
  ce29440

• Moved partial sort to own file

  

  NB4444

  authored and

  Beanavil

  committed
  736c7db

• Added partial_sort_config

  

  NB4444

  authored and

  Beanavil

  committed
  e2460c4

• Merge with nth_element_remove_oracle branch

  

  NB4444

  authored and

  Beanavil

  committed
  82cb298

• Created c++17 test for partial_sort

  

  NB4444

  authored and

  Beanavil

  committed
  dabd835

• Cleanup code based on nth_element review

  

  NB4444

  authored and

  Beanavil

  committed
  0aabab0

• Review adaptations

  

  NB4444

  authored and

  Beanavil

  committed
  ff1ed5d

• Added benchmark for partial_sort

  

  NB4444

  authored and

  Beanavil

  committed
  48e92b0

• Fixed bug with inplicit casting in partial sort

  

  NB4444

  authored and

  Beanavil

  committed
  8f68114

• add static_cast to fix compiler warning

  

  parbenc

  authored and

  Beanavil

  committed
  5e156b2

• Restored tests for device histogram_even for half/bfloat16 types

  

  Melirius

  authored and

  Beanavil

  committed
  4d51ed1

• Removed unused variable and formatting

  

  NB4444

  authored and

  Beanavil

  committed
  f5d8432

• ci: Enable debug builds excluding test_block_adjacent_difference/discontinuity

  Show description for c450a74
  

  ex-rzr

  authored and

  Beanavil

  committed
  c450a74

• test(test_device_batch_memcpy.cpp): fix invalid calls being made to generate_random_data_n

  

  Naraenda

  authored and

  Beanavil

  committed
  0cec9f8

• test(test_device_batch_memcpy.cpp): standardize test names

  

  Naraenda

  authored and

  Beanavil

  committed
  04122d2

• test(test_intrinsics.cpp): fix invalid calls being made to test_utils::get_random_data

  

  Naraenda

  authored and

  Beanavil

  committed
  cb98d2b

• ci(.gitlab-ci.yml): add hardened libc++ assertions when building tests with gitlab ci

  

  Naraenda

  authored and

  Beanavil

  committed
  78720b5

• docs: update changelog

  

  Naraenda

  authored and

  Beanavil

  committed
  62857dc

• docs: fix doxygen errors and warnings

  

  Naraenda

  authored and

  Beanavil

  committed
  1677174

• build(cmake/Dependencies.cmake): build rocm-cmake depedency during populate step when fetching it by source

  

  Naraenda

  authored and

  Beanavil

  committed
  285369a

• refactor(benchmark_config_dispatch.cpp): fix unused variable and function

  

  Naraenda

  authored and

  Beanavil

  committed
  951ea9d

• chore: bump version to 3.3.0

  

  Naraenda

  authored and

  Beanavil

  committed
  099afd3

• Reduce items_per_thread for merge_sort to one for large types

  

  NB4444

  authored and

  Beanavil

  committed
  eaddb40

• Reduce block_size for device_merge with large types

  

  NB4444

  authored and

  Beanavil

  committed
  87abd8c