Develop stream 2023 06 19 updated (#383)

* Resolve "Implement curandGetDirectionVectors32 and curandGetDirectionVectors64"

* add missing ctest test files to ci artifacts

* remove hiprand

* review comments round 1

* Fix rocRAND CUDA build instructions

* Document how to build rocRAND with a non-standard ROCm installation path

* Add documentation on passing the module path when building for NVIDIA

* Apply rule of 5 in benchmarks

* Make mt19937_engine movable

* Apply rule of 5 to rocrand.hpp

This could be a breaking change for users, but any code that no longer
compiles was probably buggy anyway.

* Apply rule of 5 / 0 to rocrand internal types

* Abort if deallocation fails in random number engines

Abort instead of silently continuing when `hipFree` fails in rng
destructors. `hipFree` should only fail if
1. there is something wrong
with the runtime in which case continuing doesn't really make sense or
2. There is a bug in rocRAND trying to double free or free something
that was not allocated with `hipMalloc`.

This is usually a sign that something went pretty wrong, and probably
other weird failures would follow if we tried to continue anyway.

This makes a test for move construction / assignment in the next commit
sound, so that it does not need to depend on implicitly terminating when
device side memory is used after freeing it.

* Various fixes in rocrand.hpp

- Add missing order setter for rocrand_cpp::threefry4x64_20_engine
- Fix typo in alias name for the mrg31k3p engine in rocrand_cpp
- Fix default ordering for rocrand_cpp::lfsr113_engine

* Test move operations for rocrand_cpp generators

* Add copyright check script based on hipCUB

* Implement device-side initialization of MT19937

* Optimize jump ahead for MT19937

* Replace % n with a simpler wrapping code;
* Change the innermost loop so i < n must be checked only for
  the last iteration, others can be unrolled without if;
* Fix block size so blockDim.x can be replaced with a constant.

* Move MT19937's jump-ahead state from shared memory to registers

* Update years in copyright

* Add a test for MT19937's jump-ahead on device

It compares the new device implementation with the old host one.
mt19937_engine is used only for the test so it has been moved from
the library to tests.

* Move precomputed jumps to a separate compilation unit

* Add polynomials for 256 ^ 2 generators

There are 510 (== (256 - 1) * 2) polynomials in the table instead of
13 polynomials of power-of-two jumps.
Each generator requires up to 2 iterations of Horner algorithm.
Added a compressed file clist_mt19937.txt created using jump_ahead_1.02

* add workaround for sdma oversubscription

* Add continuity tests for MT19937

* Add tests for heads and tails

* Optimize MT19937 and add continuity support

* Dynamic indexing of state's mt[] is removed;
* Long sequences are generated mostly with fast unrolled loops;
* Short sequences can be generated by a special kernel when possible;

* Optimize MT19937 for small sizes

mt19937_octo_engine is stored as uints.
The kernel for short sequences loads (coalesced access) values much
faster because they are stored consequentially now.

* Improve comments with implementation details

* Add tests for distribution changing

* Combine engine loading/saving routines into accessor class

* remove reliance on manually-specified WIN32 define

* Remove references to HCC

* copyright fixes

---------

Co-authored-by: Nara Prasetya <[email protected]>
Co-authored-by: Robin Voetter <[email protected]>
Co-authored-by: Nol Moonen <[email protected]>
Co-authored-by: Gergely Meszaros <[email protected]>
Co-authored-by: Anton Gorenko <[email protected]>
Co-authored-by: Bálint Soproni <[email protected]>
Co-authored-by: Beatriz Navidad Vilches <[email protected]>

Author: 8 people

Diff for: .clang-format

@@ -41,7 +41,7 @@ AllowShortLoopsOnASingleLine: false
 AlwaysBreakAfterReturnType: None
 AlwaysBreakBeforeMultilineStrings: false
 AlwaysBreakTemplateDeclarations: Yes
-AttributeMacros: ['HIPRANDAPI', 'QUALIFIERS', 'FQUALIFIERS']
+AttributeMacros: ['QUALIFIERS', 'FQUALIFIERS']
 BinPackArguments: false
 BinPackParameters: false
 BitFieldColonSpacing: Both

Diff for: .githooks/pre-commit

@@ -3,8 +3,27 @@
 # Redirect output to stderr.
 exec 1>&2
 
+check_failed=false
+
 # Do the code format check
 if ! "$(git rev-parse --show-toplevel)/scripts/code-format/check-format.sh" HEAD --cached 1>&2; then
+    printf "\n\033[31mFailed\033[0m: code format check.\n"
+    check_failed=true
+fi
+
+# Do the copyright check
+# update & apply copyright when hook config is set, otherwise just verify
+opts="-qc"
+if [ "$(git config --get --type bool --default false hooks.updateCopyright)" = "true" ]; then
+    opts="-qca"
+fi
+
+if ! "$(git rev-parse --show-toplevel)/scripts/copyright-date/check-copyright.sh" "$opts" 1>&2; then
+    printf "\n\033[31mFailed\033[0m: copyright date check.\n"
+    check_failed=true
+fi
+
+if $check_failed; then
     printf "
 Pre-commit check failed, please fix the reported errors.
 Note: Use '\033[33mgit commit --no-verify\033[0m' to bypass checks.\n"

Diff for: .gitlab-ci.yml

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2023 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
@@ -60,6 +60,20 @@ clang-format:
     - git config --global --add safe.directory $CI_PROJECT_DIR
     - scripts/code-format/check-format.sh $CI_MERGE_REQUEST_DIFF_BASE_SHA --binary "$CLANG_FORMAT"
 
+copyright-date:
+  extends:
+    - .deps:rocm
+  stage: lint
+  needs: []
+  tags:
+    - rocm-build
+  rules:
+    - if: '$CI_PIPELINE_SOURCE == "merge_request_event"'
+  script:
+    - cd $CI_PROJECT_DIR
+    - git config --global --add safe.directory $CI_PROJECT_DIR
+    - scripts/copyright-date/check-copyright.sh -v -d $CI_MERGE_REQUEST_DIFF_BASE_SHA
+
 .rocm:cmake-minimum:
   variables:
     COMPILER: hipcc
@@ -198,7 +212,6 @@ clang-format:
       -D BUILD_FILE_REORG_BACKWARD_COMPATIBILITY=OFF
       -D BUILD_SHARED_LIBS=${BUILD_SHARED_LIBS}
       -D USE_HIP_CPU=ON
-      -D BUILD_HIPRAND=OFF
       -D CMAKE_INSTALL_PREFIX=${ROCRAND_DIR}
     - cmake --build $CI_PROJECT_DIR/build
     - cmake --build $CI_PROJECT_DIR/build --target package
@@ -210,6 +223,7 @@ clang-format:
       - $CI_PROJECT_DIR/build/library/
       - $CI_PROJECT_DIR/build/test/test_*
       - $CI_PROJECT_DIR/build/test/CTestTestfile.cmake
+      - $CI_PROJECT_DIR/build/test/linkage/CTestTestfile.cmake
       - $CI_PROJECT_DIR/build/benchmark/benchmark_*
       - $CI_PROJECT_DIR/build/deps/googlebenchmark/
       - $CI_PROJECT_DIR/build/deps/hip-cpu/
@@ -331,7 +345,10 @@ build-cpu:cmake-latest:
   stage: test
   script:
     - cd $CI_PROJECT_DIR/build
-    - ctest --output-on-failure
+    # Parallel execution (with other AMDGPU processes) can oversubscribe the SDMA queue.
+    # This causes the hipMemcpy to fail, which is not reported as an error by HIP.
+    # As a temporary workaround, disable the SDMA for test stability.
+    - HSA_ENABLE_SDMA=0 ctest --output-on-failure
 
 test:rocm:
   needs:
@@ -462,7 +479,6 @@ test:nvcc-python:
       -B $CI_PROJECT_DIR/build_package_test
       -D CMAKE_CXX_COMPILER=${COMPILER}
       -D USE_HIP_CPU=ON
-      -D BUILD_HIPRAND=OFF
       -D TBB_DIR=$CI_PROJECT_DIR/build/tbb-src/cmake
       -D hip_cpu_rt_DIR=$CI_PROJECT_DIR/build/deps/hip-cpu/share/hip_cpu_rt/cmake
       -D rocrand_DIR=${ROCRAND_DIR}
@@ -532,7 +548,6 @@ test:hip-cpu-package:
       -D BUILD_FILE_REORG_BACKWARD_COMPATIBILITY=OFF
       -D CMAKE_CXX_COMPILER=${COMPILER}
       -D USE_HIP_CPU=ON
-      -D BUILD_HIPRAND=OFF
       -D DEPENDENCIES_FORCE_DOWNLOAD=ON
       -D CMAKE_INSTALL_PREFIX=${ROCRAND_DIR}
     - $SUDO_CMD cmake --build $CI_PROJECT_DIR/build_only_install --target install

Diff for: .gitmodules

@@ -1,12 +1,42 @@
 # Change Log for rocRAND
 
 Full documentation for rocRAND is available at [https://rocrand.readthedocs.io/en/latest/](https://rocrand.readthedocs.io/en/latest/)
+
+## (Unreleased) rocRAND-x.x.x for ROCm 6.0.0
+### Changed
+- Removed hipRAND submodule from rocRAND. hipRAND is now only available as a separate package.
+- Generator classes from `rocrand.hpp` are no longer copyable, in previous versions these copies
+would copy internal references to the generators and would lead to double free or memory leak errors.
+  These types should be moved instead of copied, and move constructors and operators are now defined
+  for them. 
+- Improved MT19937 initialization and generation performance.
+- Removed references to and workarounds for deprecated hcc
+
+### Fixed
+- `mt19937_engine` from `rocrand.hpp` is now move-constructible and move-assignable. Previously the
+move constructor and move assignment operator was deleted for this class.
+- Various fixes for the C++ wrapper header rocrand.hpp
+  - fixed the name of `mrg31k3p` it is now correctly spelled (was incorrectly named`mrg31k3a` in
+    previous versions).
+  - added missing `order` setter method for `threefry4x64`
+  - fixed the default ordering parameter for `lfsr113`
+
 ## (Unreleased) rocRAND-2.10.17 for ROCm 5.5.0
 ### Added
 - MT19937 pseudo random number generator based on M. Matsumoto and T. Nishimura, 1998, Mersenne Twister: A 623-dimensionally equidistributed uniform pseudorandom number generator.
 - New benchmark for the device API using Google Benchmark, `benchmark_rocrand_device_api`, replacing `benchmark_rocrand_kernel`. `benchmark_rocrand_kernel` is deprecated and will be removed in a future version. Likewise, `benchmark_curand_host_api` is added to replace `benchmark_curand_generate` and `benchmark_curand_device_api` is added to replace `benchmark_curand_kernel`.
 - experimental HIP-CPU feature
 - ThreeFry pseudorandom number generator based on Salmon et al., 2011, "Parallel random numbers: as easy as 1, 2, 3".
+- Accessor methods for sobol 32 and 64 direction vectors and constants:
+  - Enum `rocrand_direction_vector_set` to select the direction vector set.
+  - `rocrand_get_direction_vectors32(...)` supersedes:
+    - `rocrand_h_sobol32_direction_vectors` 
+    - `rocrand_h_scrambled_sobol32_direction_vectors`
+  - `rocrand_get_direction_vectors64(...)` supersedes:
+    - `rocrand_h_sobol64_direction_vectors` 
+    - `rocrand_h_scrambled_sobol64_direction_vectors`
+  - `rocrand_get_scramble_constants32(...)` supersedes `h_scrambled_sobol32_constants`
+  - `rocrand_get_scramble_constants64(...)` supersedes `h_scrambled_sobol64_constants`
 ### Changed
 - Python 2.7 is no longer officially supported.
 

Diff for: CHANGELOG.md

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2018-2022 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2018-2023 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
@@ -104,7 +104,7 @@ if (AMDGPU_TARGETS)
     set( AMDGPU_TARGETS "${gpus}" CACHE STRING "AMD GPU targets to compile for" FORCE )
 endif()
 
-# Verify that hcc compiler is used on ROCM platform
+# Verify that hipcc compiler is used on ROCM platform
 # TODO: Fix VerifyCompiler for Windows
 if(NOT USE_HIP_CPU)
   if(NOT WIN32)
@@ -173,9 +173,7 @@ print_configuration_summary()
 # Tools
 add_subdirectory(tools)
 
-# rocRAND and hipRAND libraries
-# This adds library targets: rocrand, hiprand,
-# also includes Fortran wrapper
+# rocRAND library, adds library target "rocrand" and includes Fortran wrapper
 add_subdirectory(library)
 
 if(BUILD_TEST OR BUILD_BENCHMARK)
@@ -229,23 +227,6 @@ endif()
 
 set(ROCRAND_CONFIG_DIR "\${CPACK_PACKAGING_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR}" CACHE PATH "Path placed into ldconfig file")
 
-if(WIN32)
-    option(BUILD_HIPRAND "Build hipRAND as part of rocRAND" OFF)
-else()
-    option(BUILD_HIPRAND "Build hipRAND as part of rocRAND" ON)
-endif()
-if(BUILD_HIPRAND)
-    if(HIP_COMPILER STREQUAL "nvcc")
-        set(BUILD_WITH_LIB "CUDA")
-    endif()
-    if(EXISTS ${CMAKE_SOURCE_DIR}/hipRAND/CMakeLists.txt)
-        set(HIPRAND_SUBMODULE ON CACHE BOOL "Build hipRAND as a submodule of rocRAND." FORCE)
-        add_subdirectory(${CMAKE_SOURCE_DIR}/hipRAND ${CMAKE_BINARY_DIR}/hipRAND)
-    else()
-        message(FATAL_ERROR "BUILD_HIPRAND was set, but the hipRAND submodule could not be found. Use git submodule update --init to clone the hipRAND submodule, or set BUILD_HIPRAND to OFF (--no-hiprand flag for ./install)")
-    endif()
-endif()
-
 rocm_create_package(
     NAME ${package_name}
     DESCRIPTION "Radeon Open Compute RAND library"

Diff for: CMakeLists.txt

@@ -9,7 +9,7 @@ programming language and optimised for AMD's latest discrete GPUs. It is designe
 of AMD's Radeon Open Compute [ROCm](https://rocm.github.io/) runtime, but it also works on
 CUDA enabled GPUs.
 
-Prior to ROCm version 5.0, this project included the [hipRAND](https://github.com/ROCmSoftwarePlatform/hipRAND.git) wrapper. As of version 5.0, this has been split into a separate library.
+Prior to ROCm version 5.0, this project included the [hipRAND](https://github.com/ROCmSoftwarePlatform/hipRAND.git) wrapper. As of version 5.0, this has been split into a separate library. As of version 6.0, hipRAND can no longer be built from rocRAND.
 
 ## Supported Random Number Generators
 
@@ -89,23 +89,26 @@ cd rocRAND; mkdir build; cd build
 
 # Configure rocRAND, setup options for your system
 # Build options: BUILD_TEST (off by default), BUILD_BENCHMARK (off by default), BUILD_SHARED_LIBS (on by default)
+# Additionally, the ROCm installation prefix should be passed using CMAKE_PREFIX_PATH or by setting the ROCM_PATH environment variable.
 #
 # ! IMPORTANT !
 # Set C++ compiler to HIP-clang. You can do it by adding 'CXX=<path-to-compiler>'
 # before 'cmake' or setting cmake option 'CMAKE_CXX_COMPILER' to path to the compiler.
 #
 # The python interface do not work with static library.
 #
-[CXX=hipcc] cmake -DBUILD_BENCHMARK=ON ../. # or cmake-gui ../.
+[CXX=hipcc] cmake -DBUILD_BENCHMARK=ON ../. -DCMAKE_PREFIX_PATH=/opt/rocm # or cmake-gui ../.
 
-# To configure rocRAND for Nvidia platforms, 'CXX=<path-to-nvcc>', `CXX=nvcc` or omitting the flag
-# entirely before 'cmake' is sufficient
-[CXX=nvcc] cmake -DBUILD_BENCHMARK=ON ../. # or cmake-gui ../.
+# To configure rocRAND for NVIDIA platforms, the CXX compiler must be set to a host compiler. The CUDA compiler can
+# be set explicitly using `-DCMAKE_CUDA_COMPILER=<path-to-nvcc>`.
+# Additionally, the path to FindHIP.cmake should be passed via CMAKE_MODULE_PATH. By default, this is module is
+# installed in /opt/rocm/hip/cmake.
+cmake -DBUILD_BENCHMARK=ON ../. -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_MODULE_PATH=/opt/rocm/hip/cmake # or cmake-gui ../.
 # or
-cmake -DBUILD_BENCHMARK=ON ../. # or cmake-gui ../.
+[CXX=g++] cmake -DBUILD_BENCHMARK=ON -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_MODULE_PATH=/opt/rocm/hip/cmake ../. # or cmake-gui ../.
 
-# To configure rocRAND for HIP-CPU (experimental), the USE_HIP_CPU flag is required and BUILD_HIPRAND should be turned off
-[CXX=g++] cmake -DUSE_HIP_CPU=ON -DBUILD_HIPRAND=OFF -DBUILD_BENCHMARK=ON ../. # or cmake-gui ../.
+# To configure rocRAND for HIP-CPU (experimental), the USE_HIP_CPU flag is required
+[CXX=g++] cmake -DUSE_HIP_CPU=ON -DBUILD_BENCHMARK=ON -DCMAKE_PREFIX_PATH=/opt/rocm ../. # or cmake-gui ../.
 
 # Build
 make -j4
@@ -231,10 +234,9 @@ cd rocRAND; cd build
 
 ## Wrappers
 
-* C++ wrappers for host API of rocRAND and hipRAND are in files [`rocrand.hpp`](./library/include/rocrand/rocrand.hpp)
-and [`hiprand.hpp`](./library/include/rocrand/hiprand.hpp).
+* C++ wrappers for host API of rocRAND are in [`rocrand.hpp`](./library/include/rocrand/rocrand.hpp).
 * [Fortran wrappers](./library/src/fortran/).
-* [Python wrappers](./python/): [rocRAND](./python/rocrand) and [hipRAND](./python/hiprand).
+* [Python wrappers](./python/): [rocRAND](./python/rocrand).
 
 ## Support
 

Diff for: README.md

@@ -1,4 +1,4 @@
-// Copyright (c) 2017-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2023 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
@@ -117,6 +117,11 @@ struct runner
         CUDA_CALL(cudaDeviceSynchronize());
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));
@@ -190,6 +195,11 @@ struct runner<curandStateMtgp32_t>
         CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));
@@ -297,6 +307,11 @@ struct runner<curandStateSobol32_t>
         CUDA_CALL(cudaFree(directions));
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));
@@ -369,6 +384,11 @@ struct runner<curandStateScrambledSobol32_t>
         CUDA_CALL(cudaFree(scramble_constants));
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));
@@ -425,6 +445,11 @@ struct runner<curandStateSobol64_t>
         CUDA_CALL(cudaFree(directions));
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));
@@ -496,6 +521,11 @@ struct runner<curandStateScrambledSobol64_t>
         CUDA_CALL(cudaFree(scramble_constants));
     }
 
+    runner(const runner&)            = delete;
+    runner(runner&&)                 = delete;
+    runner& operator=(const runner&) = delete;
+    runner& operator=(runner&&)      = delete;
+
     ~runner()
     {
         CUDA_CALL(cudaFree(states));