Develop stream 2024-07-01

.clang-format

@@ -45,7 +45,6 @@ AllowShortLoopsOnASingleLine: false
 AlwaysBreakAfterReturnType: None
 AlwaysBreakBeforeMultilineStrings: false
 AlwaysBreakTemplateDeclarations: Yes
-AttributeMacros: ['QUALIFIERS', 'FQUALIFIERS']
 BinPackArguments: false
 BinPackParameters: false
 BitFieldColonSpacing: Both
@@ -88,7 +87,7 @@ EmptyLineAfterAccessModifier: Never
 EmptyLineBeforeAccessModifier: Always
 ExperimentalAutoDetectBinPacking: false
 FixNamespaceComments: true
-ForEachMacros: [ foreach, Q_FOREACH, BOOST_FOREACH ]
+ForEachMacros: []
 IfMacros: []
 IncludeBlocks: Preserve
 IndentAccessModifiers: false
@@ -139,4 +138,17 @@ SpacesInConditionalStatement: false
 SpacesInContainerLiterals: true
 SpacesInParentheses: false
 SpacesInSquareBrackets: false
+
+AttributeMacros: ['__host__', '__device__', '__global__', '__forceinline__', '__shared__', '__launch_bounds__']
+# Trick clang into thinking that our C-style attributes are C++-style attributes
+Macros:
+    - __host__=[[host]]
+    - __device__=[[device]]
+    - __global__=[[global]]
+    - __forceinline__=[[forceinline]]
+    - __shared__=[[shared]]
+    - __launch_bounds__(x)=[[launch_bounds(x)]]
+    - __attribute__(x)=[[attribute(x)]]
+BreakAfterAttributes: Always
+
 ---

.gitignore

@@ -88,3 +88,4 @@ CMakeLists.txt.user
 
 # Python
 __pycache__
+.venv

.gitlab-ci.yml

@@ -22,7 +22,6 @@
 
 variables:
   GIT_SUBMODULE_STRATEGY: normal # Fetch submodules before job start (non-recursively)
-  ROCM_PATH: /opt/rocm
 
 include:
   - project: amd/ci-templates
@@ -31,6 +30,7 @@ include:
       - /defaults.yaml
       - /deps-cmake.yaml
       - /deps-docs.yaml
+      - /deps-format.yaml
       - /deps-rocm.yaml
       - /deps-nvcc.yaml
       - /deps-windows.yaml
@@ -46,20 +46,7 @@ stages:
 
 clang-format:
   extends:
-    - .deps:rocm
-  stage: lint
-  needs: []
-  tags:
-    - build
-  variables:
-    CLANG_FORMAT: "/opt/rocm/llvm/bin/clang-format"
-    GIT_CLANG_FORMAT: "/opt/rocm/llvm/bin/git-clang-format"
-  rules:
-    - if: '$CI_PIPELINE_SOURCE == "merge_request_event"'
-  script:
-    - cd $CI_PROJECT_DIR
-    - git config --global --add safe.directory $CI_PROJECT_DIR
-    - scripts/code-format/check-format.sh $CI_MERGE_REQUEST_DIFF_BASE_SHA --binary "$CLANG_FORMAT"
+    - .lint:clang-format
 
 copyright-date:
   extends:
@@ -144,12 +131,10 @@ copyright-date:
     - .rules:build
   needs: []
   script:
-    - $SUDO_CMD apt-get install -y flang
     - cmake
       -S $CI_PROJECT_DIR
       -B $CI_PROJECT_DIR/build
       -G Ninja
-      -D CMAKE_FORTRAN_COMPILER=/usr/bin/flang
       -D CMAKE_CXX_COMPILER=${COMPILER}
       -D BUILD_TEST=ON
       -D BUILD_BENCHMARK=ON
@@ -173,12 +158,10 @@ copyright-date:
     - .rules:build
   needs: []
   script:
-    - $SUDO_CMD apt-get install -y flang
     - cmake
       -S $CI_PROJECT_DIR
       -B $CI_PROJECT_DIR/build
       -G Ninja
-      -D CMAKE_FORTRAN_COMPILER=/usr/bin/flang
       -D CMAKE_CXX_COMPILER=${COMPILER}
       -D CMAKE_CUDA_HOST_COMPILER=${COMPILER}
       -D BUILD_TEST=ON
@@ -353,7 +336,6 @@ benchmark:benchmark-tuning-generate-results:
     - .rules:test
   stage: test
   script:
-    - $SUDO_CMD apt-get install -y flang
     - cd $CI_PROJECT_DIR/build
     # 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.
@@ -598,10 +580,9 @@ test:nvcc-parity:
       -D BUILD_SHARED_LIBS="$BUILD_SHARED_LIBS"
       -D BUILD_TEST=ON
       -D CMAKE_BUILD_TYPE=Release
-      -D CMAKE_CXX_COMPILER:FILEPATH="$ROCM_PATH/bin/clang++.exe"
+      -D CMAKE_CXX_COMPILER:FILEPATH="${env:HIP_PATH}/bin/clang++.exe"
       -D CMAKE_INSTALL_PREFIX:PATH="$CI_PROJECT_DIR/build/install"
-      -D CMAKE_PREFIX_PATH:PATH="$ROCM_PATH/lib/cmake"
-      -D DEPENDENCIES_FORCE_DOWNLOAD=ON
+      -D CMAKE_PREFIX_PATH:PATH="${env:HIP_PATH}/lib/cmake"
       -D DISABLE_WERROR=OFF *>&1
     # Building
     - cmake --build "$CI_PROJECT_DIR/build" *>&1
@@ -629,8 +610,8 @@ test:windows:
       -B "$CI_PROJECT_DIR/build_install_test"
       -G Ninja
       -D CMAKE_BUILD_TYPE=Release
-      -D CMAKE_CXX_COMPILER="$ROCM_PATH/bin/clang++.exe"
-      -D CMAKE_PREFIX_PATH="$ROCM_PATH/lib/cmake;$CI_PROJECT_DIR/build/install" *>&1
+      -D CMAKE_CXX_COMPILER:FILEPATH="${env:HIP_PATH}/bin/clang++.exe"
+      -D CMAKE_PREFIX_PATH:FILEPATH="${env:HIP_PATH}/lib/cmake;$CI_PROJECT_DIR/build/install" *>&1
     # Build package test
     - cmake --build "$CI_PROJECT_DIR/build_install_test"
     # Copy rocRAND.dll to the package test build directory

CHANGELOG.md

@@ -3,11 +3,24 @@
 Documentation for rocRAND is available at
 [https://rocm.docs.amd.com/projects/rocRAND/en/latest/](https://rocm.docs.amd.com/projects/rocRAND/en/latest/)
 
+## (Unreleased) rocRAND-3.2.0 for ROCm 6.3.0
+
+### Additions
+
+* Added host generator for MT19937
+* Support for `rocrand_generate_poisson` in hipGraphs
+
+### Changes
+
+* `rocrand_discrete` for MTGP32, LFSR113 and ThreeFry generators now uses the alias method, which is faster than binary search in CDF.
+
 ## (Unreleased) rocRAND-3.1.1 for ROCm 6.2.0
 
 ## Fixes
-  * Fixed " unknown extension ?>" issue in scripts/config-tuning/select_best_config.py 
-    when using python version thats older than 3.11
+
+* Fixed " unknown extension ?>" issue in scripts/config-tuning/select_best_config.py 
+  when using python version thats older than 3.11
+* Fixed low random sequence quality of `ROCRAND_RNG_PSEUDO_THREEFRY2_64_20` and `ROCRAND_RNG_PSEUDO_THREEFRY4_64_20`.
 
 ## (Unreleased) rocRAND-3.1.0 for ROCm 6.2.0
 

CMakeLists.txt

@@ -33,6 +33,7 @@ cmake_dependent_option(BUILD_BENCHMARK_TUNING
 option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
 option(CODE_COVERAGE "Build with code coverage flags (clang only)" OFF)
 option(DEPENDENCIES_FORCE_DOWNLOAD "Don't search the system for dependencies, always download them" OFF)
+cmake_dependent_option(RUN_SLOW_TESTS "Run extra tests with CTest. These cover niche functionality and take long time" OFF "BUILD_TEST" OFF)
 
 # Install prefix
 set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "Install path prefix, prepended onto install directories")

README.md

@@ -235,6 +235,4 @@ Bugs and feature requests can be reported through the
 Contributions of any kind are most welcome! You can find more information at
 [CONTRIBUTING](./CONTRIBUTING.md).
 
-Licensing information is located at [LICENSE](./LICENSE.txt). Note that [statistical tests](./test/crush) link
-to the TestU01 library distributed under GNU General Public License (GPL) version 3. Therefore, the GPL
-version 3 license applies to that part of the project.
+Licensing information is located at [LICENSE](./LICENSE.txt).

benchmark/benchmark_rocrand_host_api.cpp

@@ -179,6 +179,8 @@ int main(int argc, char* argv[])
     std::vector<rng_type_t> benchmarked_engine_types{ROCRAND_RNG_PSEUDO_LFSR113,
                                                      ROCRAND_RNG_PSEUDO_MRG31K3P,
                                                      ROCRAND_RNG_PSEUDO_MRG32K3A,
+                                                     ROCRAND_RNG_PSEUDO_MTGP32,
+                                                     ROCRAND_RNG_PSEUDO_MT19937,
                                                      ROCRAND_RNG_PSEUDO_PHILOX4_32_10,
                                                      ROCRAND_RNG_PSEUDO_THREEFRY2_32_20,
                                                      ROCRAND_RNG_PSEUDO_THREEFRY2_64_20,
@@ -190,12 +192,6 @@ int main(int argc, char* argv[])
                                                      ROCRAND_RNG_QUASI_SOBOL64,
                                                      ROCRAND_RNG_QUASI_SCRAMBLED_SOBOL64};
 
-    if(!benchmark_host)
-    {
-        benchmarked_engine_types.push_back(ROCRAND_RNG_PSEUDO_MTGP32);
-        benchmarked_engine_types.push_back(ROCRAND_RNG_PSEUDO_MT19937);
-    }
-
     const std::map<rocrand_ordering, std::string> ordering_name_map{
         {ROCRAND_ORDERING_PSEUDO_DEFAULT, "default"},
         { ROCRAND_ORDERING_PSEUDO_LEGACY,  "legacy"},

benchmark/tuning/benchmark_tuning.hpp

@@ -88,7 +88,10 @@ void run_benchmark(benchmark::State& state, const benchmark_config& config)
     generator.set_stream(stream);
 
     const auto generate_func = [&]
-    { return generator.generate(data, size, default_distribution<Distribution>{}(config)); };
+    {
+        default_distribution<Distribution> default_distribution_provider;
+        return generator.generate(data, size, default_distribution_provider(config));
+    };
 
     // Warm-up
     ROCRAND_CHECK(generate_func());
@@ -147,9 +150,7 @@ class generator_benchmark_factory
             if constexpr(std::is_same_v<T, unsigned int>)
             {
                 // The poisson distribution is only supported for unsigned int.
-                using poisson_distribution_t = rocrand_impl::host::poisson_distribution<
-                    rocrand_impl::host::DISCRETE_METHOD_ALIAS>;
-                add_benchmarks_impl<T, poisson_distribution_t>();
+                add_benchmarks_impl<T, select_poisson_distribution_t<GeneratorTemplate>>();
             }
         }
         else if constexpr(std::is_floating_point_v<T> || std::is_same_v<T, half>)

benchmark/tuning/benchmarked_generators.hpp

@@ -107,7 +107,9 @@ using mtgp32_generator_template
                                                     ConfigProvider>;
 
 template<class ConfigProvider>
-using mt19937_generator_template = rocrand_impl::host::mt19937_generator_template<ConfigProvider>;
+using mt19937_generator_template
+    = rocrand_impl::host::mt19937_generator_template<rocrand_impl::system::device_system,
+                                                     ConfigProvider>;
 
 template<class ConfigProvider>
 using philox4x32_10_generator_template

benchmark/tuning/distribution_traits.hpp

@@ -146,6 +146,15 @@ struct distribution_name<
     }
 };
 
+template<>
+struct distribution_name<rocrand_impl::host::mrg_poisson_distribution>
+{
+    std::string operator()()
+    {
+        return "poisson_unsigned_int";
+    }
+};
+
 template<class Distribution>
 struct default_distribution
 {
@@ -191,11 +200,49 @@ struct default_distribution<
 {
     auto operator()(const benchmark_config& config)
     {
-        return rocrand_impl::host::poisson_distribution<rocrand_impl::host::DISCRETE_METHOD_ALIAS>(
-            config.lambda);
+        return std::get<
+            rocrand_impl::host::poisson_distribution<rocrand_impl::host::DISCRETE_METHOD_ALIAS>>(
+            m_poisson_manager.get_distribution(config.lambda));
+    }
+
+private:
+    rocrand_impl::host::poisson_distribution_manager<rocrand_impl::host::DISCRETE_METHOD_ALIAS>
+        m_poisson_manager;
+};
+
+template<>
+struct default_distribution<rocrand_impl::host::mrg_poisson_distribution>
+{
+    auto operator()(const benchmark_config& config)
+    {
+        auto poisson_distribution = std::get<
+            rocrand_impl::host::poisson_distribution<rocrand_impl::host::DISCRETE_METHOD_ALIAS>>(
+            m_poisson_manager.get_distribution(config.lambda));
+        return rocrand_impl::host::mrg_poisson_distribution(poisson_distribution);
     }
+
+private:
+    rocrand_impl::host::poisson_distribution_manager<rocrand_impl::host::DISCRETE_METHOD_ALIAS>
+        m_poisson_manager;
 };
 
+template<template<class> class GeneratorTemplate>
+struct select_poisson_distribution
+{
+    using dummy_generator_t = GeneratorTemplate<rocrand_impl::host::static_config_provider<0, 0>>;
+    static constexpr inline rocrand_rng_type rng_type = dummy_generator_t::type();
+    static constexpr inline bool             is_mrg
+        = rng_type == ROCRAND_RNG_PSEUDO_MRG31K3P || rng_type == ROCRAND_RNG_PSEUDO_MRG32K3A;
+
+    using type = std::conditional_t<
+        is_mrg,
+        rocrand_impl::host::mrg_poisson_distribution,
+        rocrand_impl::host::poisson_distribution<rocrand_impl::host::DISCRETE_METHOD_ALIAS>>;
+};
+
+template<template<class> class GeneratorTemplate>
+using select_poisson_distribution_t = typename select_poisson_distribution<GeneratorTemplate>::type;
+
 } // namespace benchmark_tuning
 
 #endif // ROCRAND_BENCHMARK_TUNING_DISTRIBUTION_TRAITS_HPP_

docs/data-type-support.rst → docs/conceptual/data-type-support.rst

@@ -1,3 +1,9 @@
+.. meta::
+   :description: rocRAND documentation and API reference library
+   :keywords: rocRAND, ROCm, API, documentation, cuRAND
+
+.. _data-type-support:
+
 Data type support
 ******************************************
 

docs/conceptual/dynamic_ordering_configuration.rst

@@ -22,14 +22,14 @@ Building the tuning benchmarks
 
 The principle of the tuning is very simple: the random number generation kernel is run for a list of kernel block size / kernel grid size combinations, and the fastest combination is selected as the dynamic ordering configuration for the particular device. rocRAND provides an executable target that runs the benchmarks with all these combinations: `benchmark_rocrand_tuning`. This target is disabled by default, and can be enabled and built by the following snippet.
 
-Use the `AMDGPU_TARGET` variable to specify the comma-separated list of GPU architectures to build the benchmarks for. To acquire the architecture of the GPU(s) installed, run `rocminfo`, and look for `gfx` in the "ISA Info" section. ::
+Use the `GPU_TARGETS` variable to specify the comma-separated list of GPU architectures to build the benchmarks for. To acquire the architecture of the GPU(s) installed, run `rocminfo`, and look for `gfx` in the "ISA Info" section. ::
 
     $ cd rocRAND
     $ cmake -S . -B ./build
         -D BUILD_BENCHMARK=ON
         -D BUILD_BENCHMARK_TUNING=ON
         -D CMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++
-        -D AMDGPU_TARGETS=gfx908
+        -D GPU_TARGETS=gfx908
     $ cmake --build build --target benchmark_rocrand_tuning
 
 Additionally, the following CMake cache variables control the generation of the benchmarked matrix: