Merge pull request #77 from ROCmSoftwarePlatform/develop

Develop to Master for 2.6

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.6.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)

Jenkinsfile

@@ -49,10 +49,16 @@ rocprimCI:
         platform, project->
 
         def command = """
-                      echo "rocPRIM is a header only library and does not need packaging"
+                      set -x
+                      cd ${project.paths.project_build_prefix}/build
+                      make package
+                      rm -rf package && mkdir -p package
+                      mv *.deb package/
+                      dpkg -c package/*.deb
                       """
 
         platform.runCommand(this, command)
+        platform.archiveArtifacts(this, """${project.paths.project_build_prefix}/build/package/*.deb""")
     }
 
     buildProject(rocprim, formatCheck, nodes.dockerArray, compileCommand, testCommand, packageCommand)

install

@@ -15,6 +15,7 @@ function display_help()
     #    echo "    [-d|--dependencies] install build dependencies"
     echo "    [-c|--clients] build library clients too (combines with -i & -d)"
     echo "    [-g|--debug] -DCMAKE_BUILD_TYPE=Debug (default is =Release)"
+    echo "    [--hip-clang] build library for amdgpu backend using hip-clang"
 }
 
 
@@ -25,6 +26,7 @@ install_package=false
 build_package=false
 build_clients=false
 build_release=true
+build_hip_clang=false
 run_tests=false
 rocm_path=/opt/rocm/bin
 # #################################################
@@ -34,7 +36,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,package --options hicdtp -- "$@")
+    GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,debug,hip-clang,test,package --options hicdtp -- "$@")
 else
     echo "Need a new version of getopt"
     exit 1
@@ -69,6 +71,9 @@ while true; do
 	-t|--test)
 	    run_tests=true
 	    shift ;;
+	--hip-clang)
+	    build_hip_clang=true
+	    shift ;;
 	--) shift ; break ;;
 	*)  echo "Unexpected command line parameter received; aborting";
 	    exit 1
@@ -100,8 +105,12 @@ fi
 # before 'cmake' or setting cmake option 'CMAKE_CXX_COMPILER' to path to the HCC compiler.
 #
 
+compiler="hcc"
+if [[ "${build_hip_clang}" == true ]]; then
+    compiler="hipcc"
+fi
 
-CXX=$rocm_path/hcc cmake -DBUILD_BENCHMARK=ON ../../. # or cmake-gui ../.
+CXX=$rocm_path/${compiler} cmake -DBUILD_BENCHMARK=ON ../../. # or cmake-gui ../.
 
 # Build
 make -j$(nproc)

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}