Local memory limit fix at scan, reduce and select

CMakeLists.txt

@@ -73,15 +73,15 @@ include(cmake/Dependencies.cmake)
 set(AMDGPU_TARGETS gfx803;gfx900;gfx906 CACHE STRING "List of specific machine types for library to target")
 
 # Setup VERSION
-set ( VERSION_STRING "2.5.0." )
+set(VERSION_STRING "2.5.0.")
 
-# Check if BUILD_NUMBER is defined in a Jenkins envirnment
-if ($ENV{BUILD_NUMBER} )
+# Check if BUILD_NUMBER is defined in a Jenkins environment
+if($ENV{BUILD_NUMBER})
   string(CONCAT BUILD_VERSION ${VERSION_STRING} $ENV{BUILD_NUMBER})
-else ()
+else()
   string(CONCAT BUILD_VERSION ${VERSION_STRING} "0")
-endif ()
-rocm_setup_version( VERSION ${BUILD_VERSION} NO_GIT_TAG_VERSION )
+endif()
+rocm_setup_version(VERSION ${BUILD_VERSION} NO_GIT_TAG_VERSION)
 
 # Print configuration summary
 include(cmake/Summary.cmake)

install

@@ -10,6 +10,7 @@ function display_help()
     echo "./install [-h|--help] "
     echo "    [-h|--help] prints this help message"
     echo "    [-i|--install] install after build"
+    echo "    [-p]--package build package"
     #Not implemented yet
     #    echo "    [-d|--dependencies] install build dependencies"
     echo "    [-c|--clients] build library clients too (combines with -i & -d)"
@@ -21,6 +22,7 @@ function display_help()
 # global variables
 # #################################################
 install_package=false
+build_package=false
 build_clients=false
 build_release=true
 run_tests=false
@@ -32,7 +34,7 @@ rocm_path=/opt/rocm/bin
 # check if we have a modern version of getopt that can handle whitespace and long parameters
 getopt -T
 if [[ $? -eq 4 ]]; then
-    GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,debug,test --options hicdt -- "$@")
+    GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,debug,test,package --options hicdtp -- "$@")
 else
     echo "Need a new version of getopt"
     exit 1
@@ -55,7 +57,10 @@ while true; do
 	-i|--install)
 	    install_package=true
 	    shift ;;
-	-c|--clients)
+	-p|--package)
+	    build_package=true
+	    shift ;;
+	 -c|--clients)
 	    build_clients=true
 	    shift ;;
 	-g|--debug)
@@ -99,7 +104,7 @@ fi
 CXX=$rocm_path/hcc cmake -DBUILD_BENCHMARK=ON ../../. # or cmake-gui ../.
 
 # Build
-make -j32
+make -j$(nproc)
 
 if ($run_tests); then
 # Optionally, run tests if they're enabled.
@@ -110,3 +115,8 @@ if ($install_package); then
 # Install
 make install
 fi
+
+if ($build_package); then
+# Install
+make package -j$(nproc)
+fi

rocprim/CMakeLists.txt

@@ -53,9 +53,6 @@ set(INCLUDE_INSTALL_DIR ${CMAKE_INSTALL_INCLUDEDIR})
 # does not support header-only libraries (INTERFACE targets)
 rocm_install_targets(
   TARGETS rocprim rocprim_hip
-#   INCLUDE
-#   ${CMAKE_SOURCE_DIR}/rocprim/include
-#   ${CMAKE_BINARY_DIR}/rocprim/include
   PREFIX rocprim
 )
 
@@ -72,7 +69,7 @@ install(
 
 # Export targets
 rocm_export_targets(
-  TARGETS rocprim-targets
+  TARGETS roc::rocprim roc::rocprim_hip
   PREFIX rocprim
   NAMESPACE roc::
 )

rocprim/include/rocprim/device/device_reduce_by_key_config.hpp

@@ -65,7 +65,7 @@ struct reduce_by_key_config_803
             (sizeof(Key) <= 8 && sizeof(Value) <= 8),
             reduce_by_key_config<scan, kernel_config<256, 7> >
         >,
-        reduce_by_key_config<scan, kernel_config<256, ::rocprim::max(1u, 15u / item_scale)> >
+        reduce_by_key_config<scan, kernel_config<limit_block_size<256U, sizeof(Key) + sizeof(Value)>::value, ::rocprim::max(1u, 15u / item_scale)> >
     >;
 };
 
@@ -82,7 +82,7 @@ struct reduce_by_key_config_900
             (sizeof(Key) <= 8 && sizeof(Value) <= 8),
             reduce_by_key_config<scan, kernel_config<256, 10> >
         >,
-        reduce_by_key_config<scan, kernel_config<256, ::rocprim::max(1u, 15u / item_scale)> >
+        reduce_by_key_config<scan, kernel_config<limit_block_size<256U, sizeof(Key) + sizeof(Value)>::value, ::rocprim::max(1u, 15u / item_scale)> >
     >;
 };
 

rocprim/include/rocprim/device/device_reduce_config.hpp

@@ -65,7 +65,7 @@ struct reduce_config_803
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = reduce_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 16u / item_scale),
         ::rocprim::block_reduce_algorithm::using_warp_reduce
     >;
@@ -78,7 +78,7 @@ struct reduce_config_900
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = reduce_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 16u / item_scale),
         ::rocprim::block_reduce_algorithm::using_warp_reduce
     >;

rocprim/include/rocprim/device/device_scan_config.hpp

@@ -79,7 +79,7 @@ struct scan_config_803
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = scan_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 16u / item_scale),
         ROCPRIM_DETAIL_USE_LOOKBACK_SCAN,
         ::rocprim::block_load_method::block_load_transpose,
@@ -95,7 +95,7 @@ struct scan_config_900
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = scan_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 16u / item_scale),
         ROCPRIM_DETAIL_USE_LOOKBACK_SCAN,
         ::rocprim::block_load_method::block_load_transpose,

rocprim/include/rocprim/device/device_select_config.hpp

@@ -74,7 +74,7 @@ struct select_config_803
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = select_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 13u / item_scale),
         ::rocprim::block_load_method::block_load_transpose,
         ::rocprim::block_load_method::block_load_transpose,
@@ -89,7 +89,7 @@ struct select_config_900
         ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
 
     using type = select_config<
-        256,
+        limit_block_size<256U, sizeof(Value)>::value,
         ::rocprim::max(1u, 15u / item_scale),
         ::rocprim::block_load_method::block_load_transpose,
         ::rocprim::block_load_method::block_load_transpose,

test/extra/CMakeLists.txt

@@ -57,7 +57,7 @@ function(add_rocprim_test TEST_NAME TEST_SOURCES)
   add_executable(${TEST_TARGET} ${TEST_SOURCES})
   target_link_libraries(${TEST_TARGET}
     PRIVATE
-      roc::rocprim_hip
+      ${rocprim_LIBRARIES} # roc::rocprim_hip
   )
   foreach(amdgpu_target ${AMDGPU_TARGETS})
     target_link_libraries(${TEST_TARGET}