From 2e05322b290bac7046e33575f9e66ef4b040a284 Mon Sep 17 00:00:00 2001 From: Matteo Concas Date: Mon, 11 Sep 2023 11:42:14 +0200 Subject: [PATCH] Move hipification into dedicated files --- CMakeLists.txt | 2 + .../ITS/tracking/GPU/hip/CMakeLists.txt | 52 ++++++------- GPU/GPUTracking/Base/cuda/CMakeLists.txt | 6 +- GPU/GPUbenchmark/Shared/Utils.h | 5 ++ GPU/GPUbenchmark/cuda/Kernels.cu | 6 -- GPU/GPUbenchmark/hip/CMakeLists.txt | 45 ++--------- cmake/O2AddHipifiedExecutable.cmake | 76 +++++++++++++++++++ cmake/O2AddHipifiedLibrary.cmake | 71 +++++++++++++++++ 8 files changed, 187 insertions(+), 76 deletions(-) create mode 100644 cmake/O2AddHipifiedExecutable.cmake create mode 100644 cmake/O2AddHipifiedLibrary.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index ff7cf309dfa7a..49e11831095d9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -93,6 +93,8 @@ include(O2DataFile) include(O2TargetManPage) include(O2AddWorkflow) include(O2SetROOTPCMDependencies) +include(O2AddHipifiedExecutable) +include (O2AddHipifiedLibrary) # Main targets of the project in various subdirectories. Order matters. add_subdirectory(Common) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt b/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt index 8fccb7f2a87d9..6203b701e2caa 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt +++ b/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt @@ -10,34 +10,34 @@ # or submit itself to any jurisdiction. if(HIP_ENABLED) - # Hipify-perl to generate HIP sources - set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl") - file(GLOB CUDA_SOURCES_FULL_PATH "../cuda/*.cu") - foreach(file ${CUDA_SOURCES_FULL_PATH}) - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file}) - get_filename_component(CUDA_SOURCE ${file} NAME) - string(REPLACE ".cu" "" CUDA_SOURCE_NAME ${CUDA_SOURCE}) - add_custom_command( - OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip - COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} - ) - endforeach() + # # Hipify-perl to generate HIP sources + # set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl") + # file(GLOB CUDA_SOURCES_FULL_PATH "../cuda/*.cu") + # foreach(file ${CUDA_SOURCES_FULL_PATH}) + # set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file}) + # get_filename_component(CUDA_SOURCE ${file} NAME) + # string(REPLACE ".cu" "" CUDA_SOURCE_NAME ${CUDA_SOURCE}) + # add_custom_command( + # OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip + # COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip + # DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} + # ) + # endforeach() - install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/${baseTargetName} - DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) + # install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/${baseTargetName} + # DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} ${O2_HIP_CMAKE_CXX_FLAGS} -fgpu-rdc") message(STATUS "Building ITS HIP tracker") - o2_add_library(ITStrackingHIP - SOURCES ClusterLinesGPU.hip - Context.hip - TimeFrameGPU.hip - Stream.hip - TrackerTraitsGPU.hip - TracerGPU.hip - VertexerTraitsGPU.hip - Utils.hip + o2_add_hipified_library(ITStrackingHIP + SOURCES ../cuda/ClusterLinesGPU.cu + ../cuda/Context.cu + ../cuda/TimeFrameGPU.cu + ../cuda/Stream.cu + ../cuda/TrackerTraitsGPU.cu + ../cuda/TracerGPU.cu + ../cuda/VertexerTraitsGPU.cu + ../cuda/Utils.cu PUBLIC_INCLUDE_DIRECTORIES ../ PUBLIC_LINK_LIBRARIES O2::ITStracking hip::host @@ -45,8 +45,8 @@ if(HIP_ENABLED) hip::hipcub TARGETVARNAME targetName) - target_compile_definitions( - ${targetName} PRIVATE $) + # target_compile_definitions( + # ${targetName} PRIVATE $) if(O2_HIP_CMAKE_LINK_FLAGS) # Need to add gpu target also to link flags due to gpu-rdc option diff --git a/GPU/GPUTracking/Base/cuda/CMakeLists.txt b/GPU/GPUTracking/Base/cuda/CMakeLists.txt index ed719de239103..f933f7e8c1822 100644 --- a/GPU/GPUTracking/Base/cuda/CMakeLists.txt +++ b/GPU/GPUTracking/Base/cuda/CMakeLists.txt @@ -50,7 +50,7 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") set(RTC_CUDA_ARCH "750") endif() separate_arguments(CUDARTC_FLAGS) - + # convenience variables if(ALIGPU_BUILD_TYPE STREQUAL "Standalone") get_filename_component(GPUDIR ${CMAKE_SOURCE_DIR}/../ ABSOLUTE) @@ -78,7 +78,7 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") VERBATIM ) create_binary_resource(${CURTC_BIN}.command ${CURTC_BIN}.command.o) - + set(SRCS ${SRCS} ${CURTC_BIN}.src.o ${CURTC_BIN}.command.o) endif() # -------------------------------- End RTC ------------------------------------------------------- @@ -124,8 +124,6 @@ if(ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") install(TARGETS ${targetName} ARCHIVE DESTINATION lib LIBRARY DESTINATION lib) install(FILES ${HDRS} DESTINATION include) - - endif() if(ALIGPU_BUILD_TYPE STREQUAL "Standalone") diff --git a/GPU/GPUbenchmark/Shared/Utils.h b/GPU/GPUbenchmark/Shared/Utils.h index cf4e07baaae77..97d6788eb01de 100644 --- a/GPU/GPUbenchmark/Shared/Utils.h +++ b/GPU/GPUbenchmark/Shared/Utils.h @@ -47,6 +47,11 @@ exit(EXIT_FAILURE); #endif +template +void discardResult(const T&) +{ +} + enum class Test { Read, Write, diff --git a/GPU/GPUbenchmark/cuda/Kernels.cu b/GPU/GPUbenchmark/cuda/Kernels.cu index 49494531e8d16..c455e8f6126f4 100644 --- a/GPU/GPUbenchmark/cuda/Kernels.cu +++ b/GPU/GPUbenchmark/cuda/Kernels.cu @@ -428,7 +428,6 @@ float GPUbenchmark::runSequential(void (*kernel)(chunk_t*, size_t, T... // Warm up (*kernel)<<>>(chunkPtr, getBufferCapacity(chunk.second, mOptions.prime), args...); - cudaDeviceSynchronize(); GPUCHECK(cudaGetLastError()); GPUCHECK(cudaEventCreate(&start)); GPUCHECK(cudaEventCreate(&stop)); @@ -436,7 +435,6 @@ float GPUbenchmark::runSequential(void (*kernel)(chunk_t*, size_t, T... GPUCHECK(cudaEventRecord(start)); for (auto iLaunch{0}; iLaunch < nLaunches; ++iLaunch) { // Schedule all the requested kernel launches (*kernel)<<>>(chunkPtr, getBufferCapacity(chunk.second, mOptions.prime), args...); // NOLINT: clang-tidy false-positive - cudaDeviceSynchronize(); GPUCHECK(cudaGetLastError()); } GPUCHECK(cudaEventRecord(stop)); // record checkpoint @@ -772,7 +770,6 @@ void GPUbenchmark::runTest(Test test, Mode mode, KernelConfig config) } else { std::cout << "" << measurement << "\t" << iChunk << "\t" << throughput << "\t" << chunkSize << "\t" << result << std::endl; } - } } else if (mode == Mode::Concurrent) { if (!mOptions.raw) { @@ -807,7 +804,6 @@ void GPUbenchmark::runTest(Test test, Mode mode, KernelConfig config) } else { std::cout << "" << measurement << "\t" << iChunk << "\t" << throughput << "\t" << chunkSize << "\t" << results[iChunk] << std::endl; } - } if (mState.testChunks.size() > 1) { if (!mOptions.raw) { @@ -852,9 +848,7 @@ void GPUbenchmark::runTest(Test test, Mode mode, KernelConfig config) } else { std::cout << "" << measurement << "\t" << 0 << "\t" << throughput << "\t" << tot << "\t" << result << std::endl; } - } - } } diff --git a/GPU/GPUbenchmark/hip/CMakeLists.txt b/GPU/GPUbenchmark/hip/CMakeLists.txt index 4cbdf52221570..1459464d37548 100644 --- a/GPU/GPUbenchmark/hip/CMakeLists.txt +++ b/GPU/GPUbenchmark/hip/CMakeLists.txt @@ -10,43 +10,8 @@ # or submit itself to any jurisdiction. message(STATUS "Building GPU HIP benchmark") -# Hipify-perl to generate HIP sources -set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl") -file(GLOB CUDA_SOURCES_FULL_PATH "../cuda/*.cu") -foreach(file ${CUDA_SOURCES_FULL_PATH}) - set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file}) - get_filename_component(CUDA_SOURCE ${file} NAME) - string(REPLACE ".cu" "" CUDA_SOURCE_NAME ${CUDA_SOURCE}) - add_custom_command( - OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip - COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE} - ) -endforeach() -message(STATUS "CMAKE_EXE_LINKER_FLAGS: ${CMAKE_EXE_LINKER_FLAGS}") -message(STATUS "CMAKE_SHARED_LINKER_FLAGS: ${CMAKE_SHARED_LINKER_FLAGS}") -message(STATUS "CMAKE_HIP_COMPILE_FEATURES: ${CMAKE_HIP_COMPILE_FEATURES}") -message(STATUS "CMAKE_LINKER: ${CMAKE_LINKER}") -message(STATUS "CMAKE_HIP_STANDARD_LIBRARIES: ${CMAKE_HIP_STANDARD_LIBRARIES}") -message(STATUS "CMAKE_HIP_LINK_EXECUTABLE: ${CMAKE_HIP_LINK_EXECUTABLE}") -message(STATUS "CMAKE_HIP_LINK_FLAGS: ${CMAKE_HIP_LINK_FLAGS}") -message(STATUS "CMAKE_HIP_RUNTIME_LIBRARIES_STATIC: ${CMAKE_HIP_RUNTIME_LIBRARIES_STATIC}") -message(STATUS "CMAKE_HIP_RUNTIME_LIBRARIES_SHARED: ${CMAKE_HIP_RUNTIME_LIBRARIES_SHARED}") -message(STATUS "CMAKE_MODULE_LINKER_FLAGS: ${CMAKE_MODULE_LINKER_FLAGS}") -o2_add_executable(gpu-memory-benchmark-hip - SOURCES benchmark.hip - Kernels.hip - PUBLIC_LINK_LIBRARIES hip::host - hip-lang::device - Boost::program_options - TARGETVARNAME targetName) -# >:-) -include_directories(${Boost_INCLUDE_DIRS}) -add_executable(alternative-gpu-memory-benchmark-hip benchmark.hip Kernels.hip) -target_link_libraries(alternative-gpu-memory-benchmark-hip Boost::program_options ROOT::Tree) -install(TARGETS alternative-gpu-memory-benchmark-hip RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}) - -# if(O2_HIP_CMAKE_LINK_FLAGS) - # Need to add gpu target also to link flags due to gpu-rdc option - # target_link_options(${targetName} PUBLIC ${O2_HIP_CMAKE_LINK_FLAGS}) -# endif() +o2_add_hipified_executable(gpu-memory-benchmark-hip + SOURCES ../cuda/benchmark.cu + ../cuda/Kernels.cu + PUBLIC_LINK_LIBRARIES hip::host + Boost::program_options) \ No newline at end of file diff --git a/cmake/O2AddHipifiedExecutable.cmake b/cmake/O2AddHipifiedExecutable.cmake new file mode 100644 index 0000000000000..5938f493317a5 --- /dev/null +++ b/cmake/O2AddHipifiedExecutable.cmake @@ -0,0 +1,76 @@ +# Copyright 2019-2020 CERN and copyright holders of ALICE O2. +# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +# All rights not expressly granted are reserved. +# +# This software is distributed under the terms of the GNU General Public +# License v3 (GPL Version 3), copied verbatim in the file "COPYING". +# +# In applying this license CERN does not waive the privileges and immunities +# granted to it by virtue of its status as an Intergovernmental Organization +# or submit itself to any jurisdiction. + +include_guard() + +include(O2AddExecutable) + +function(o2_add_hipified_executable baseTargetName) + # Parse arguments in the same way o2_add_executable does + cmake_parse_arguments(PARSE_ARGV + 1 + A + "IS_TEST;NO_INSTALL;IS_BENCHMARK" + "COMPONENT_NAME;TARGETVARNAME" + "SOURCES;PUBLIC_LINK_LIBRARIES;JOB_POOL") + + # Process each .cu file to generate a .hip file + set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl") + set(HIP_SOURCES) + + foreach(file ${A_SOURCES}) + get_filename_component(ABS_CUDA_SORUCE ${file} ABSOLUTE) + if(file MATCHES "\\.cu$") + set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file}) + get_filename_component(CUDA_SOURCE ${file} NAME) + string(REPLACE ".cu" ".hip" HIP_SOURCE ${CUDA_SOURCE}) + set(OUTPUT_HIP_FILE "${CMAKE_CURRENT_SOURCE_DIR}/${HIP_SOURCE}") + list(APPEND HIP_SOURCES ${OUTPUT_HIP_FILE}) + + add_custom_command( + OUTPUT ${OUTPUT_HIP_FILE} + COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${ABS_CUDA_SORUCE} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${OUTPUT_HIP_FILE} + DEPENDS ${file} + ) + else() + list(APPEND HIP_SOURCES ${file}) + endif() + endforeach() + + # This is a bit cumbersome, but it seems the only suitable since cmake_parse_arguments is not capable to filter only the SOURCE variadic values + set(FORWARD_ARGS "") + if(A_IS_TEST) + list(APPEND FORWARD_ARGS "IS_TEST") + endif() + if(A_NO_INSTALL) + list(APPEND FORWARD_ARGS "NO_INSTALL") + endif() + if(A_IS_BENCHMARK) + list(APPEND FORWARD_ARGS "IS_BENCHMARK") + endif() + if(A_COMPONENT_NAME) + list(APPEND FORWARD_ARGS "COMPONENT_NAME" ${A_COMPONENT_NAME}) + endif() + if(A_TARGETVARNAME) + list(APPEND FORWARD_ARGS "TARGETVARNAME" ${A_TARGETVARNAME}) + endif() + if(A_PUBLIC_LINK_LIBRARIES) + list(APPEND FORWARD_ARGS "PUBLIC_LINK_LIBRARIES" ${A_PUBLIC_LINK_LIBRARIES}) + endif() + if(A_JOB_POOL) + list(APPEND FORWARD_ARGS "JOB_POOL" ${A_JOB_POOL}) + endif() + + # Call o2_add_executable with new sources + o2_add_executable("${baseTargetName}" + SOURCES ${HIP_SOURCES} + ${FORWARD_ARGS}) +endfunction() diff --git a/cmake/O2AddHipifiedLibrary.cmake b/cmake/O2AddHipifiedLibrary.cmake new file mode 100644 index 0000000000000..96e0a73575272 --- /dev/null +++ b/cmake/O2AddHipifiedLibrary.cmake @@ -0,0 +1,71 @@ +# Copyright 2019-2020 CERN and copyright holders of ALICE O2. +# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +# All rights not expressly granted are reserved. +# +# This software is distributed under the terms of the GNU General Public +# License v3 (GPL Version 3), copied verbatim in the file "COPYING". +# +# In applying this license CERN does not waive the privileges and immunities +# granted to it by virtue of its status as an Intergovernmental Organization +# or submit itself to any jurisdiction. + +include_guard() + +include(O2AddLibrary) + +function(o2_add_hipified_library baseTargetName) + # Parse arguments in the same way o2_add_library does + cmake_parse_arguments(PARSE_ARGV + 1 + A + "" + "TARGETVARNAME" + "SOURCES;PUBLIC_INCLUDE_DIRECTORIES;PUBLIC_LINK_LIBRARIES;PRIVATE_INCLUDE_DIRECTORIES;PRIVATE_LINK_LIBRARIES" + ) + + # Process each .cu file to generate a .hip file + set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl") + set(HIP_SOURCES) + + foreach(file ${A_SOURCES}) + get_filename_component(ABS_CUDA_SORUCE ${file} ABSOLUTE) + if(file MATCHES "\\.cu$") + set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file}) + get_filename_component(CUDA_SOURCE ${file} NAME) + string(REPLACE ".cu" ".hip" HIP_SOURCE ${CUDA_SOURCE}) + set(OUTPUT_HIP_FILE "${CMAKE_CURRENT_SOURCE_DIR}/${HIP_SOURCE}") + list(APPEND HIP_SOURCES ${OUTPUT_HIP_FILE}) + + add_custom_command( + OUTPUT ${OUTPUT_HIP_FILE} + COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${ABS_CUDA_SORUCE} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${OUTPUT_HIP_FILE} + DEPENDS ${file} + ) + else() + list(APPEND HIP_SOURCES ${file}) + endif() + endforeach() + + # This is a bit cumbersome, but it seems the only suitable since cmake_parse_arguments is not capable to filter only the SOURCE variadic values + set(FORWARD_ARGS "") + if(A_TARGETVARNAME) + list(APPEND FORWARD_ARGS "TARGETVARNAME" ${A_TARGETVARNAME}) + endif() + if(A_PUBLIC_INCLUDE_DIRECTORIES) + list(APPEND FORWARD_ARGS "PUBLIC_INCLUDE_DIRECTORIES" ${A_PUBLIC_INCLUDE_DIRECTORIES}) + endif() + if(A_PUBLIC_LINK_LIBRARIES) + list(APPEND FORWARD_ARGS "PUBLIC_LINK_LIBRARIES" ${A_PUBLIC_LINK_LIBRARIES}) + endif() + if(A_PRIVATE_INCLUDE_DIRECTORIES) + list(APPEND FORWARD_ARGS "PRIVATE_INCLUDE_DIRECTORIES" ${A_PRIVATE_INCLUDE_DIRECTORIES}) + endif() + if(A_PRIVATE_LINK_LIBRARIES) + list(APPEND FORWARD_ARGS "PRIVATE_LINK_LIBRARIES" ${A_PRIVATE_LINK_LIBRARIES}) + endif() + + # Call o2_add_library with new sources + o2_add_library("${baseTargetName}" + SOURCES ${HIP_SOURCES} + ${FORWARD_ARGS}) +endfunction() \ No newline at end of file