From 1d9cbc6a6f658bbcb6079648ef8b28e896460cb6 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 26 Nov 2024 19:55:34 +0100 Subject: [PATCH 1/3] GPU CMake: Add check that RTC source files do not contain system headers --- GPU/GPUTracking/Base/cuda/CMakeLists.txt | 8 +++++++- GPU/GPUTracking/Base/hip/CMakeLists.txt | 17 +++++++++++++---- GPU/GPUTracking/Base/opencl2/CMakeLists.txt | 1 + 3 files changed, 21 insertions(+), 5 deletions(-) diff --git a/GPU/GPUTracking/Base/cuda/CMakeLists.txt b/GPU/GPUTracking/Base/cuda/CMakeLists.txt index a24092f50ebaf..995b9224a4ad0 100644 --- a/GPU/GPUTracking/Base/cuda/CMakeLists.txt +++ b/GPU/GPUTracking/Base/cuda/CMakeLists.txt @@ -69,7 +69,7 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") add_custom_command( OUTPUT ${GPU_RTC_BIN}.src COMMAND cat ${GPUDIR}/Base/cuda/GPUReconstructionCUDAIncludes.h > ${GPU_RTC_BIN}.src - COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_CUDA_STANDARD} -D__CUDA_ARCH__=${RTC_CUDA_ARCH} -D__CUDACC__ -x c++ -E ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src + COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_CUDA_STANDARD} -D__CUDA_ARCH__=${RTC_CUDA_ARCH} -D__CUDACC__ -x c++ -nostdinc -E ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src MAIN_DEPENDENCY ${GPU_RTC_SRC} IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC} COMMAND_EXPAND_LISTS @@ -77,6 +77,12 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") ) create_binary_resource(${GPU_RTC_BIN}.src ${GPU_RTC_BIN}.src.o) + add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_CUDA_SRC_CHK.done + COMMAND ! grep "# [0-9]* \"\\(/usr/\\|.*GCC-Toolchain\\)" ${GPU_RTC_BIN}.src > ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_CUDA_SRC_CHK.done || bash -c "echo ERROR: CUDA RTC sources contain standard headers 1>&2 && exit 1" + COMMENT Checking CUDA RTC File ${GPU_RTC_BIN}.src + DEPENDS ${GPU_RTC_BIN}.src VERBATIM) + add_custom_target(${MODULE}_CUDA_SRC_CHK ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_CUDA_SRC_CHK.done) + add_custom_command( OUTPUT ${GPU_RTC_BIN}.command COMMAND echo -n "${CMAKE_CUDA_COMPILER} ${GPU_RTC_FLAGS_SEPARATED} ${GPU_RTC_DEFINES} -fatbin" > ${GPU_RTC_BIN}.command diff --git a/GPU/GPUTracking/Base/hip/CMakeLists.txt b/GPU/GPUTracking/Base/hip/CMakeLists.txt index f488ce8c7dd14..40b095143d639 100644 --- a/GPU/GPUTracking/Base/hip/CMakeLists.txt +++ b/GPU/GPUTracking/Base/hip/CMakeLists.txt @@ -52,8 +52,11 @@ if(NOT DEFINED GPUCA_HIP_HIPIFY_FROM_CUDA OR "${GPUCA_HIP_HIPIFY_FROM_CUDA}") list(APPEND HIP_SOURCES "${GPUCA_HIP_SOURCE_DIR}/${HIP_SOURCE}") endforeach() - add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done COMMAND diff -u ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip ${CMAKE_CURRENT_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip && touch ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done DEPENDS ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip ${CMAKE_CURRENT_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip) - add_custom_target(${MODULE}_HIPIFIED_CHK DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done) + add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done + COMMAND diff -u ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip ${CMAKE_CURRENT_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip > ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done + DEPENDS ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip ${CMAKE_CURRENT_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip + COMMENT Checking HIPified file ${CMAKE_CURRENT_SOURCE_DIR}/GPUReconstructionHIPkernel.template.hip) + add_custom_target(${MODULE}_HIPIFIED_CHK ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIPIFIED_CHK.done) else() get_filename_component(GPUCA_HIP_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR} ABSOLUTE) endif() @@ -103,7 +106,7 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") add_custom_command( OUTPUT ${GPU_RTC_BIN}.src COMMAND cat ${GPUDIR}/Base/hip/GPUReconstructionHIPIncludes.h > ${GPU_RTC_BIN}.src - COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_HIP_STANDARD} -D__HIPCC__ -D__HIP_DEVICE_COMPILE__ -x c++ -E ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src + COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_HIP_STANDARD} -D__HIPCC__ -D__HIP_DEVICE_COMPILE__ -x c++ -nostdinc -E ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src MAIN_DEPENDENCY ${GPU_RTC_SRC} IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC} COMMAND_EXPAND_LISTS @@ -111,6 +114,12 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") ) create_binary_resource(${GPU_RTC_BIN}.src ${GPU_RTC_BIN}.src.o) + add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIP_SRC_CHK.done + COMMAND ! grep "# [0-9]* \"\\(/usr/\\|.*GCC-Toolchain\\)" ${GPU_RTC_BIN}.src > ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIP_SRC_CHK.done || bash -c "echo ERROR: HIP RTC sources contain standard headers 1>&2 && exit 1" + COMMENT Checking HIP RTC File ${GPU_RTC_BIN}.src + DEPENDS ${GPU_RTC_BIN}.src VERBATIM) + add_custom_target(${MODULE}_HIP_SRC_CHK ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/${MODULE}_HIP_SRC_CHK.done) + add_custom_command( OUTPUT ${GPU_RTC_BIN}.command COMMAND echo -n "${hip_HIPCC_EXECUTABLE} ${GPU_RTC_FLAGS_SEPARATED} ${GPU_RTC_DEFINES} --genco" > ${GPU_RTC_BIN}.command @@ -202,7 +211,7 @@ target_include_directories(${MODULE}_CXX PRIVATE $ ${CL_BIN}.src MAIN_DEPENDENCY ${CL_SRC} IMPLICIT_DEPENDS CXX ${CL_SRC} From 510508bbbbc913a73a2e940f001f002bb2d2c9df Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 26 Nov 2024 19:55:41 +0100 Subject: [PATCH 2/3] Remove obsolete tools --- .../Standalone/tools/rtc/rtcsource.sh | 15 ---- GPU/GPUTracking/Standalone/tools/rtc/test.cu | 86 ------------------- 2 files changed, 101 deletions(-) delete mode 100755 GPU/GPUTracking/Standalone/tools/rtc/rtcsource.sh delete mode 100644 GPU/GPUTracking/Standalone/tools/rtc/test.cu diff --git a/GPU/GPUTracking/Standalone/tools/rtc/rtcsource.sh b/GPU/GPUTracking/Standalone/tools/rtc/rtcsource.sh deleted file mode 100755 index 9f59855eaf278..0000000000000 --- a/GPU/GPUTracking/Standalone/tools/rtc/rtcsource.sh +++ /dev/null @@ -1,15 +0,0 @@ -#!/bin/bash -cat < source.cu -# 1 "/home/qon/alice/O2/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu" -# 1 "/home/qon/alice/O2/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtcPre.h" 1 -EOT -cat src/Base/cuda/GPUReconstructionCUDAIncludes.h >> source.cu -nvcc -std=c++17 -gencode arch=compute_75,code=sm_75 -E \ - -I src/ -I src/Common/ -I src/Base/ -I src/SliceTracker/ -I src/Merger/ -I src/TRDTracking/ -I src/TPCClusterFinder/ -I src/TPCConvert/ -I src/Global/ -I src/dEdx/ -I src/TPCFastTransformation/ -I src/GPUUtils/ -I src/DataCompression -I src/ITS \ - -I$HOME/alice/O2/DataFormats/Detectors/TPC/include -I$HOME/alice/O2/Detectors/Base/include -I$HOME/alice/O2/Detectors/Base/src -I$HOME/alice/O2/Common/MathUtils/include -I$HOME/alice/O2/DataFormats/Headers/include \ - -I$HOME/alice/O2/Detectors/TRD/base/include -I$HOME/alice/O2/Detectors/TRD/base/src -I$HOME/alice/O2/Detectors/ITSMFT/ITS/tracking/include -I$HOME/alice/O2/Detectors/ITSMFT/ITS/tracking/cuda/include -I$HOME/alice/O2/Common/Constants/include \ - -I$HOME/alice/O2/DataFormats/common/include -I$HOME/alice/O2/DataFormats/Detectors/TRD/include -I$HOME/alice/O2/Detectors/Raw/include \ - -DGPUCA_HAVE_O2HEADERS -DGPUCA_TPC_GEOMETRY_O2 -DGPUCA_STANDALONE \ - ~/alice/O2/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu \ -| sed '1,/^# 1 ".*GPUReconstructionCUDArtcPre.h" 1$/d' | grep -v O2_GPU_KERNEL_TEMPLATE_REPLACE \ ->> source.cu diff --git a/GPU/GPUTracking/Standalone/tools/rtc/test.cu b/GPU/GPUTracking/Standalone/tools/rtc/test.cu deleted file mode 100644 index f567218046ac7..0000000000000 --- a/GPU/GPUTracking/Standalone/tools/rtc/test.cu +++ /dev/null @@ -1,86 +0,0 @@ -#include -#include -#include - -#define NVRTC_SAFE_CALL(x) \ - do { \ - nvrtcResult result = x; \ - if (result != NVRTC_SUCCESS) { \ - std::cerr << "\nerror: " #x " failed with error " \ - << nvrtcGetErrorString(result) << '\n'; \ - exit(1); \ - } \ - } while (0) - -#define CUDA_SAFE_CALL(x) \ - do { \ - CUresult result = x; \ - if (result != CUDA_SUCCESS) { \ - const char* msg; \ - cuGetErrorName(result, &msg); \ - std::cerr << "\nerror: " #x " failed with error " \ - << msg << '\n'; \ - exit(1); \ - } \ - } while (0) - -int32_t main(int argc, char** argv) -{ - //Read Sourcecode from file - uint32_t filesize; - FILE* pFile; - //Open file - if ((pFile = fopen("source.cu", "rb")) == NULL) - exit(1); - //Optain File Size - fseek(pFile, 0, SEEK_END); - filesize = ftell(pFile); - rewind(pFile); - //Read file - char* sourceCode = new char[filesize + 1]; - if (fread(sourceCode, 1, filesize, pFile) != filesize) - exit(1); - //Make sourceCode 0-terminated - sourceCode[filesize] = 0; - fclose(pFile); - - nvrtcProgram prog; - NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, // prog - sourceCode, // buffer - "saxpy.cu", // name - 0, // numHeaders - NULL, // headers - NULL)); // includeNames - delete[] sourceCode; - //const char *opts[] = {"-default-device -std=c++17 --extended-lambda -Xptxas -O4 -Xcompiler -O4 -use_fast_math --ftz=true"}; - const char* opts[] = {"-default-device", "--std=c++17", "-use_fast_math", "-ftz=true"}; - nvrtcResult compileResult = nvrtcCompileProgram(prog, // prog - sizeof(opts) / sizeof(opts[0]), // numOptions - opts); // options - size_t logSize; - NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize)); - char* log = new char[logSize]; - NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log)); - std::cout << log << '\n'; - delete[] log; - if (compileResult != NVRTC_SUCCESS) { - exit(1); - } - size_t ptxSize; - NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize)); - char* ptx = new char[ptxSize]; - NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx)); - NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog)); - CUmodule module; - CUfunction kernel; - CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0)); - CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "foo")); - void* args[] = {}; - CUDA_SAFE_CALL( - cuLaunchKernel(kernel, - 1, 1, 1, // grid dim - 32, 1, 1, // block dim - 0, NULL, // shared mem and stream - args, 0)); // arguments - return 0; -} From 5899144c718b64a76aa2c4a8b3abc1ede81ccfc0 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 14 Jan 2025 23:34:47 +0100 Subject: [PATCH 3/3] HIP RTC CMake: Fix dependency --- GPU/GPUTracking/Base/hip/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/GPU/GPUTracking/Base/hip/CMakeLists.txt b/GPU/GPUTracking/Base/hip/CMakeLists.txt index 40b095143d639..acd87c528e8e4 100644 --- a/GPU/GPUTracking/Base/hip/CMakeLists.txt +++ b/GPU/GPUTracking/Base/hip/CMakeLists.txt @@ -109,6 +109,7 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT") COMMAND ${CMAKE_CXX_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_HIP_STANDARD} -D__HIPCC__ -D__HIP_DEVICE_COMPILE__ -x c++ -nostdinc -E ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src MAIN_DEPENDENCY ${GPU_RTC_SRC} IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC} + DEPENDS ${MODULE}_HIPIFIED COMMAND_EXPAND_LISTS COMMENT "Preparing HIP RTC source file ${GPU_RTC_BIN}.src" )