Develop stream 2025-03-10 (#622)

* Deprecate fortran interface

* Resolve "Index performance of rocRAND on MI300 and curand on H100 and H200"

* Stop using __AMDGCN_WAVEFRONT_SIZE and warpSize as compile-time constants

* log more information during configure

* docs(changelog): re-organize newly added changelog entries

* Resolve "bug found on SOBOL generator"

---------

Co-authored-by: Beatriz Navidad Vilches <[email protected]>
Co-authored-by: Anton Gorenko <[email protected]>
Co-authored-by: Borys Petrov <[email protected]>

Author: 4 people

CHANGELOG.md

@@ -12,20 +12,25 @@ Documentation for rocRAND is available at
 ### Changed
 
 * Updated several `gfx942` auto tuning parameters.
-* Deprecated C++14 and set the default target to C++17.
-* Directly accessing the (scrambled) sobol32 and sobol64 constants and direction vectors is deprecated:
+* Modified error handling and expanded the error information for the case of double-deallocation of the (scrambled) sobol32 and sobol64 constants and direction vectors.
+
+### Removed
+
+* Removed inline assembly and the `ENABLE_INLINE_ASM` CMake option. Inline assembly was used to optimizate of multiplications in the Mrg32k3a and Philox 4x32-10 generators. It is no longer needed because the current HIP compiler is able to produce code with the same or better performance.
+* Removed instances of the deprecated clang definition `__AMDGCN_WAVEFRONT_SIZE`.
+
+### Upcoming changes
+
+* Deprecated the rocRAND Fortran API in favor of hipfort.
+* Deprecated C++14 and set the default target to C++17. C++14 will be removed in the next major release.
+* Directly accessing the (scrambled) sobol32 and sobol64 constants and direction vectors is deprecated and will be removed in the next major release. For:
   * `h_scrambled_sobol32_constants`, use `rocrand_get_scramble_constants32` instead.
   * `h_scrambled_sobol64_constants`, use `rocrand_get_scramble_constants64` instead.
   * `rocrand_h_sobol32_direction_vectors`, use `rocrand_get_direction_vectors32` instead.
   * `rocrand_h_sobol64_direction_vectors`, use `rocrand_get_direction_vectors64` instead.
   * `rocrand_h_scrambled_sobol32_direction_vectors`, use `rocrand_get_direction_vectors32` instead.
   * `rocrand_h_scrambled_sobol64_direction_vectors`, use `rocrand_get_direction_vectors64` instead.
 
-### Upcoming changes
-* C++14 will be removed in the next major release.
-* Directly accessing the (scrambled) sobol32 and sobol64 constants and direction vectors will be removed in the next major release.
-
-
 ### Fixed
 
 * Fixed an issue where `mt19937.hpp` would cause kernel errors during auto tuning.

README.md

@@ -121,10 +121,6 @@ compilers) may cause a build failure; if you encounter errors with the existing
 other dependencies, you can pass the `DEPENDENCIES_FORCE_DOWNLOAD` flag to CMake, which can
 help to solve the problem.
 
-To disable inline assembly optimizations in rocRAND (for both the host library and
-the device functions provided in `rocrand_kernel.h`), set the CMake option `ENABLE_INLINE_ASM`
-to `OFF`.
-
 ## Running unit tests
 
 ```shell

benchmark/benchmark_curand_generate.cpp

@@ -150,25 +150,31 @@ void run_benchmarks(const cli::Parser& parser,
     const std::string format = parser.get<std::string>("format");
     if (distribution == "uniform-uint")
     {
-        run_benchmark<unsigned int>(
-            parser,
-            rng_type,
-            stream,
-            [](curandGenerator_t gen, unsigned int* data, size_t size)
-            { return curandGenerate(gen, data, size); },
-            distribution,
-            engine);
+        if(rng_type != CURAND_RNG_QUASI_SOBOL64 && rng_type != CURAND_RNG_QUASI_SCRAMBLED_SOBOL64)
+        {
+            run_benchmark<unsigned int>(
+                parser,
+                rng_type,
+                stream,
+                [](curandGenerator_t gen, unsigned int* data, size_t size)
+                { return curandGenerate(gen, data, size); },
+                distribution,
+                engine);
+        }
     }
     if (distribution == "uniform-long-long")
     {
-        run_benchmark<unsigned long long>(
-            parser,
-            rng_type,
-            stream,
-            [](curandGenerator_t gen, unsigned long long* data, size_t size)
-            { return curandGenerateLongLong(gen, data, size); },
-            distribution,
-            engine);
+        if(rng_type == CURAND_RNG_QUASI_SOBOL64 || rng_type == CURAND_RNG_QUASI_SCRAMBLED_SOBOL64)
+        {
+            run_benchmark<unsigned long long>(
+                parser,
+                rng_type,
+                stream,
+                [](curandGenerator_t gen, unsigned long long* data, size_t size)
+                { return curandGenerateLongLong(gen, data, size); },
+                distribution,
+                engine);
+        }
     }
     if (distribution == "uniform-float")
     {
@@ -311,6 +317,10 @@ int main(int argc, char *argv[])
     parser.set_optional<std::vector<std::string>>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str());
     parser.set_optional<std::vector<std::string>>("engine", "engine", {"philox"}, engine_desc.c_str());
     parser.set_optional<std::vector<double>>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution");
+    parser.set_optional<std::string>("format",
+                                     "format",
+                                     {"console"},
+                                     "output format: console or csv");
     parser.run_and_exit_if_error();
 
     std::vector<std::string> engines;
@@ -365,6 +375,17 @@ int main(int argc, char *argv[])
     cudaStream_t stream;
     CUDA_CALL(cudaStreamCreate(&stream));
 
+    std::string format         = parser.get<std::string>("format");
+    bool        console_output = format.compare("console") == 0 ? true : false;
+
+    if(!console_output)
+    {
+        std::cout
+            << "Engine,Distribution,Throughput,Samples,AvgTime (1 Trial),Time(all),Size,Lambda"
+            << std::endl;
+        std::cout << ",,GB/s,GSample/s,ms),ms),values," << std::endl;
+    }
+
     for (auto engine : engines)
     {
         rng_type_t rng_type = CURAND_RNG_PSEUDO_XORWOW;
@@ -392,11 +413,13 @@ int main(int argc, char *argv[])
             exit(1);
         }
 
-        std::cout << engine << ":" << std::endl;
+        if(console_output)
+            std::cout << engine << ":" << std::endl;
 
         for (auto distribution : distributions)
         {
-            std::cout << "  " << distribution << ":" << std::endl;
+            if(console_output)
+                std::cout << "  " << distribution << ":" << std::endl;
             run_benchmarks(parser, rng_type, distribution, engine, stream);
         }
         std::cout << std::endl;

benchmark/benchmark_curand_host_api.cpp

@@ -56,7 +56,7 @@ void run_benchmark(benchmark::State&     state,
 
     if(benchmark_host)
     {
-        data = new T[size];
+        data = new T[rounded_size];
         CURAND_CALL(curandCreateGeneratorHost(&generator, rng_type));
     }
     else
@@ -84,7 +84,7 @@ void run_benchmark(benchmark::State&     state,
     // Warm-up
     for(size_t i = 0; i < 15; i++)
     {
-        CURAND_CALL(generate_func(generator, data, size));
+        CURAND_CALL(generate_func(generator, data, rounded_size));
     }
     CUDA_CALL(cudaDeviceSynchronize());
 
@@ -97,7 +97,7 @@ void run_benchmark(benchmark::State&     state,
         CUDA_CALL(cudaEventRecord(start, stream));
         for(size_t i = 0; i < trials; i++)
         {
-            CURAND_CALL(generate_func(generator, data, size));
+            CURAND_CALL(generate_func(generator, data, rounded_size));
         }
         CUDA_CALL(cudaEventRecord(stop, stream));
         CUDA_CALL(cudaEventSynchronize(stop));
@@ -108,8 +108,8 @@ void run_benchmark(benchmark::State&     state,
         state.SetIterationTime(elapsed / 1000.f);
     }
 
-    state.SetBytesProcessed(trials * state.iterations() * size * sizeof(T));
-    state.SetItemsProcessed(trials * state.iterations() * size);
+    state.SetBytesProcessed(trials * state.iterations() * rounded_size * sizeof(T));
+    state.SetItemsProcessed(trials * state.iterations() * rounded_size);
 
     CUDA_CALL(cudaEventDestroy(stop));
     CUDA_CALL(cudaEventDestroy(start));
@@ -128,6 +128,10 @@ void run_benchmark(benchmark::State&     state,
 void configure_parser(cli::Parser& parser)
 {
     parser.set_optional<size_t>("size", "size", DEFAULT_RAND_N, "number of values");
+    parser.set_optional<bool>("byte-size",
+                              "byte-size",
+                              false,
+                              "--size is interpreted as the number of generated bytes");
     parser.set_optional<size_t>("trials", "trials", 20, "number of trials");
     parser.set_optional<size_t>("offset", "offset", 0, "offset of generated pseudo-random values");
     parser.set_optional<size_t>("dimensions",
@@ -232,23 +236,26 @@ int main(int argc, char* argv[])
             const std::string name_engine_prefix
                 = benchmark_name_prefix + "<" + name + "," + ordering_name_map.at(ordering) + ",";
 
-            benchmarks.emplace_back(benchmark::RegisterBenchmark(
-                (name_engine_prefix + "uniform-uint>").c_str(),
-                &run_benchmark<unsigned int>,
-                [](curandGenerator_t gen, unsigned int* data, size_t size_gen)
-                { return curandGenerate(gen, data, size_gen); },
-                size,
-                byte_size,
-                trials,
-                dimensions,
-                offset,
-                engine_type,
-                ordering,
-                benchmark_host,
-                stream));
+            if(engine_type != CURAND_RNG_QUASI_SOBOL64
+               && engine_type != CURAND_RNG_QUASI_SCRAMBLED_SOBOL64)
 
-            if(engine_type == CURAND_RNG_QUASI_SOBOL64
-               || engine_type == CURAND_RNG_QUASI_SCRAMBLED_SOBOL64)
+            {
+                benchmarks.emplace_back(benchmark::RegisterBenchmark(
+                    (name_engine_prefix + "uniform-uint>").c_str(),
+                    &run_benchmark<unsigned int>,
+                    [](curandGenerator_t gen, unsigned int* data, size_t size_gen)
+                    { return curandGenerate(gen, data, size_gen); },
+                    size,
+                    byte_size,
+                    trials,
+                    dimensions,
+                    offset,
+                    engine_type,
+                    ordering,
+                    benchmark_host,
+                    stream));
+            }
+            else
             {
                 benchmarks.emplace_back(benchmark::RegisterBenchmark(
                     (name_engine_prefix + "uniform-long-long>").c_str(),

cmake/Summary.cmake

@@ -1,4 +1,41 @@
-function (print_configuration_summary)
+function(print_configuration_summary)
+    find_package(Git)
+    if(GIT_FOUND)
+        execute_process(
+            COMMAND ${GIT_EXECUTABLE} show --format=%H --no-patch
+            WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+            OUTPUT_VARIABLE COMMIT_HASH
+            OUTPUT_STRIP_TRAILING_WHITESPACE
+        )
+        execute_process(
+            COMMAND ${GIT_EXECUTABLE} show --format=%s --no-patch
+            WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+            OUTPUT_VARIABLE COMMIT_SUBJECT
+            OUTPUT_STRIP_TRAILING_WHITESPACE
+        )
+    endif()
+
+    execute_process(
+        COMMAND ${CMAKE_CXX_COMPILER} --version
+        WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+        OUTPUT_VARIABLE CMAKE_CXX_COMPILER_VERBOSE_DETAILS
+        OUTPUT_STRIP_TRAILING_WHITESPACE
+    )
+
+    find_program(UNAME_EXECUTABLE uname)
+    if(UNAME_EXECUTABLE)
+        execute_process(
+            COMMAND ${UNAME_EXECUTABLE} -a
+            WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+            OUTPUT_VARIABLE LINUX_KERNEL_DETAILS
+            OUTPUT_STRIP_TRAILING_WHITESPACE
+        )
+    endif()  
+
+    string(REPLACE "\n" ";" CMAKE_CXX_COMPILER_VERBOSE_DETAILS "${CMAKE_CXX_COMPILER_VERBOSE_DETAILS}")
+    list(TRANSFORM CMAKE_CXX_COMPILER_VERBOSE_DETAILS PREPEND "--     ")
+    string(REPLACE ";" "\n" CMAKE_CXX_COMPILER_VERBOSE_DETAILS "${CMAKE_CXX_COMPILER_VERBOSE_DETAILS}")
+
     message(STATUS "")
     message(STATUS "******** Summary ********")
     message(STATUS "General:")
@@ -35,4 +72,15 @@ function (print_configuration_summary)
     endif()
     message(STATUS "  BUILD_ADDRESS_SANITIZER    : ${BUILD_ADDRESS_SANITIZER}")
     message(STATUS "  DEPENDENCIES_FORCE_DOWNLOAD: ${DEPENDENCIES_FORCE_DOWNLOAD}")
+    message(STATUS "")
+    message(STATUS "Detailed:")
+    message(STATUS "  C++ compiler details       : \n${CMAKE_CXX_COMPILER_VERBOSE_DETAILS}")
+if(GIT_FOUND)
+    message(STATUS "  Commit                     : ${COMMIT_HASH}")
+    message(STATUS "                               ${COMMIT_SUBJECT}")
+endif()
+if(UNAME_EXECUTABLE)
+    message(STATUS "  Unix name                  : ${LINUX_KERNEL_DETAILS}")
+endif()
+  
 endfunction()

library/CMakeLists.txt

@@ -43,20 +43,6 @@ set(rocRAND_HIP_SRCS
     "src/rocrand_mt19937_precomputed.cpp"
     "src/rocrand.cpp")
 
-# When enabled, it defines ROCRAND_ENABLE_INLINE_ASM in rocrand_version.h, which
-# turns on inline asm in rocRAND (for both compiled library and device functions).
-option(ENABLE_INLINE_ASM "Enable inline asm optimisations in rocRAND" ON)
-if(ENABLE_INLINE_ASM)
-    set(
-        rocrand_ENABLE_INLINE_ASM
-        "\n// Enables inline asm optimisations\n"
-        "#if !defined(ROCRAND_ENABLE_INLINE_ASM) && !defined(ROCRAND_DISABLE_INLINE_ASM)\n"
-        "    #define ROCRAND_ENABLE_INLINE_ASM\n"
-        "#endif"
-    )
-    string(REPLACE ";" "" rocrand_ENABLE_INLINE_ASM "${rocrand_ENABLE_INLINE_ASM}")
-endif()
-
 # Configure a header file to pass the rocRAND version
 configure_file(
     "${PROJECT_SOURCE_DIR}/library/include/rocrand/rocrand_version.h.in"

library/include/rocrand/rocrand_common.h

@@ -1,4 +1,4 @@
-// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2025 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
@@ -92,63 +92,17 @@
 namespace rocrand_device {
 namespace detail {
 
-#if ( defined(__HIP_PLATFORM_NVCC__) || \
-      defined(__gfx801__) || \
-      defined(__gfx802__) || \
-      defined(__gfx803__) || \
-      defined(__gfx810__) || \
-      defined(__gfx900__) || \
-      defined(__gfx902__) || \
-      defined(__gfx904__) || \
-      defined(__gfx906__) || \
-      defined(__gfx908__) || \
-      defined(__gfx909__) || \
-      defined(__gfx1030__) )
-  #if !defined(ROCRAND_ENABLE_INLINE_ASM)
-    #define ROCRAND_ENABLE_INLINE_ASM
-  #endif
-#else
-  #if defined(__HIP_DEVICE_COMPILE__) && defined(ROCRAND_ENABLE_INLINE_ASM)
-    #undef ROCRAND_ENABLE_INLINE_ASM
-  #endif
-#endif
-
-__forceinline__ __device__ __host__ unsigned long long
+__forceinline__ __device__ __host__
+unsigned long long
     mad_u64_u32(const unsigned int x, const unsigned int y, const unsigned long long z)
 {
-#if defined(__HIP_PLATFORM_AMD__) && defined(__HIP_DEVICE_COMPILE__) \
-    && defined(ROCRAND_ENABLE_INLINE_ASM)
-
-    #if __AMDGCN_WAVEFRONT_SIZE == 64u
-    using sgpr_t = unsigned long long;
-  #elif __AMDGCN_WAVEFRONT_SIZE == 32u
-    using sgpr_t = unsigned int;
-  #endif
-
-    unsigned long long r;
-    sgpr_t c; // carry bits, SGPR, unused
-    // x has "r" constraint. This allows to use both VGPR and SGPR
-    // (to save VGPR) as input.
-    // y and z have "v" constraints, because only one SGPR or literal
-    // can be read by the instruction.
-    asm volatile("v_mad_u64_u32 %0, %1, %2, %3, %4"
-      : "=v"(r), "=s"(c) : "r"(x), "v"(y), "v"(z)
-    );
-    return r;
-  #elif defined(__HIP_PLATFORM_NVCC__) && defined(__HIP_DEVICE_COMPILE__) \
-        && defined(ROCRAND_ENABLE_INLINE_ASM)
-
-    unsigned long long r;
-    asm("mad.wide.u32 %0, %1, %2, %3;"
-        : "=l"(r) : "r"(x), "r"(y), "l"(z)
-    );
-    return r;
-
-  #else // host code
-
     return static_cast<unsigned long long>(x) * static_cast<unsigned long long>(y) + z;
+}
 
-  #endif
+__forceinline__ __device__ __host__
+unsigned long long mul_u64_u32(const unsigned int x, const unsigned int y)
+{
+    return static_cast<unsigned long long>(x) * static_cast<unsigned long long>(y);
 }
 
 // This helps access fields of engine's internal state which