Skip to content

Commit

Permalink
Move hipification into dedicated files
Browse files Browse the repository at this point in the history
  • Loading branch information
mconcas committed Sep 11, 2023
1 parent d454f03 commit 2e05322
Show file tree
Hide file tree
Showing 8 changed files with 187 additions and 76 deletions.
2 changes: 2 additions & 0 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
52 changes: 26 additions & 26 deletions Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -10,43 +10,43 @@
# 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
hip::device
hip::hipcub
TARGETVARNAME targetName)

target_compile_definitions(
${targetName} PRIVATE $<TARGET_PROPERTY:O2::ITStracking,COMPILE_DEFINITIONS>)
# target_compile_definitions(
# ${targetName} PRIVATE $<TARGET_PROPERTY:O2::ITStracking,COMPILE_DEFINITIONS>)

if(O2_HIP_CMAKE_LINK_FLAGS)
# Need to add gpu target also to link flags due to gpu-rdc option
Expand Down
6 changes: 2 additions & 4 deletions GPU/GPUTracking/Base/cuda/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -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 -------------------------------------------------------
Expand Down Expand Up @@ -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")
Expand Down
5 changes: 5 additions & 0 deletions GPU/GPUbenchmark/Shared/Utils.h
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,11 @@
exit(EXIT_FAILURE);
#endif

template <typename T>
void discardResult(const T&)
{
}

enum class Test {
Read,
Write,
Expand Down
6 changes: 0 additions & 6 deletions GPU/GPUbenchmark/cuda/Kernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -428,15 +428,13 @@ float GPUbenchmark<chunk_t>::runSequential(void (*kernel)(chunk_t*, size_t, T...

// Warm up
(*kernel)<<<nBlocks, nThreads, 0, stream>>>(chunkPtr, getBufferCapacity<chunk_t>(chunk.second, mOptions.prime), args...);
cudaDeviceSynchronize();
GPUCHECK(cudaGetLastError());
GPUCHECK(cudaEventCreate(&start));
GPUCHECK(cudaEventCreate(&stop));

GPUCHECK(cudaEventRecord(start));
for (auto iLaunch{0}; iLaunch < nLaunches; ++iLaunch) { // Schedule all the requested kernel launches
(*kernel)<<<nBlocks, nThreads, 0, stream>>>(chunkPtr, getBufferCapacity<chunk_t>(chunk.second, mOptions.prime), args...); // NOLINT: clang-tidy false-positive
cudaDeviceSynchronize();
GPUCHECK(cudaGetLastError());
}
GPUCHECK(cudaEventRecord(stop)); // record checkpoint
Expand Down Expand Up @@ -772,7 +770,6 @@ void GPUbenchmark<chunk_t>::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) {
Expand Down Expand Up @@ -807,7 +804,6 @@ void GPUbenchmark<chunk_t>::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) {
Expand Down Expand Up @@ -852,9 +848,7 @@ void GPUbenchmark<chunk_t>::runTest(Test test, Mode mode, KernelConfig config)
} else {
std::cout << "" << measurement << "\t" << 0 << "\t" << throughput << "\t" << tot << "\t" << result << std::endl;
}

}

}
}

Expand Down
45 changes: 5 additions & 40 deletions GPU/GPUbenchmark/hip/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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)
76 changes: 76 additions & 0 deletions cmake/O2AddHipifiedExecutable.cmake
Original file line number Diff line number Diff line change
@@ -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()
71 changes: 71 additions & 0 deletions cmake/O2AddHipifiedLibrary.cmake
Original file line number Diff line number Diff line change
@@ -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()

0 comments on commit 2e05322

Please sign in to comment.