Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly
Appearance settings

GPU: RTC and CMake improvements #13071

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidrohr merged 4 commits into from
Apr 24, 2024
Merged

GPU: RTC and CMake improvements #13071

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
4e911f0
dpl-workflow: Use /var/tmp as cache folder for GPU RTC cache
davidrohr Apr 24, 2024
47215ac
GPU: Add ignoreCacheValid debug option
davidrohr Apr 24, 2024
db76164
GPU: Fix GPU build flags once more, HIP flags got their -O3 overridde…
davidrohr Apr 24, 2024
4c515f9
GPU: Add RTC debug option to override the command line
davidrohr Apr 24, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 7 additions & 6 deletions 13 GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAGenRTC.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -93,28 +93,28 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
if (fread(sharead, 1, 20, fp) != 20) {
throw std::runtime_error("Cache file corrupt");
}
if (memcmp(sharead, shasource, 20)) {
if (!mProcessingSettings.rtc.ignoreCacheValid && memcmp(sharead, shasource, 20)) {
GPUInfo("Cache file content outdated (source)");
break;
}
if (fread(sharead, 1, 20, fp) != 20) {
throw std::runtime_error("Cache file corrupt");
}
if (memcmp(sharead, shaparam, 20)) {
if (!mProcessingSettings.rtc.ignoreCacheValid && memcmp(sharead, shaparam, 20)) {
GPUInfo("Cache file content outdated (param)");
break;
}
if (fread(sharead, 1, 20, fp) != 20) {
throw std::runtime_error("Cache file corrupt");
}
if (memcmp(sharead, shacmd, 20)) {
if (!mProcessingSettings.rtc.ignoreCacheValid && memcmp(sharead, shacmd, 20)) {
GPUInfo("Cache file content outdated (commandline)");
break;
}
if (fread(sharead, 1, 20, fp) != 20) {
throw std::runtime_error("Cache file corrupt");
}
if (memcmp(sharead, shakernels, 20)) {
if (!mProcessingSettings.rtc.ignoreCacheValid && memcmp(sharead, shakernels, 20)) {
GPUInfo("Cache file content outdated (kernel definitions)");
break;
}
Expand All @@ -123,7 +123,7 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
if (fread(&cachedSettings, sizeof(cachedSettings), 1, fp) != 1) {
throw std::runtime_error("Cache file corrupt");
}
if (memcmp(&cachedSettings, &mProcessingSettings.rtc, sizeof(cachedSettings))) {
if (!mProcessingSettings.rtc.ignoreCacheValid && memcmp(&cachedSettings, &mProcessingSettings.rtc, sizeof(cachedSettings))) {
GPUInfo("Cache file content outdated (rtc parameters)");
break;
}
Expand Down Expand Up @@ -159,6 +159,7 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
}
HighResTimer rtcTimer;
rtcTimer.ResetStart();
std::string baseCommand = getenv("O2_GPU_RTC_OVERRIDE_CMD") ? std::string(getenv("O2_GPU_RTC_OVERRIDE_CMD")) : std::string(_binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len);
#ifdef WITH_OPENMP
#pragma omp parallel for schedule(dynamic, 1)
#endif
Expand All @@ -181,7 +182,7 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
throw std::runtime_error("Error writing file");
}
fclose(fp);
std::string command = std::string(_binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len);
std::string command = baseCommand;
command += " -c " + filename + "_" + std::to_string(i) + mRtcSrcExtension + " -o " + filename + "_" + std::to_string(i) + mRtcBinExtension;
if (mProcessingSettings.debugLevel < 0) {
command += " &> /dev/null";
Expand Down
1 change: 1 addition & 0 deletions 1 GPU/GPUTracking/Definitions/GPUSettingsList.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,6 +200,7 @@ AddOption(compilePerKernel, bool, true, "", 0, "Run one RTC compilation per kern
AddOption(enable, bool, false, "", 0, "Use RTC to optimize GPU code")
AddOption(runTest, int, 0, "", 0, "Do not run the actual benchmark, but just test RTC compilation (1 full test, 2 test only compilation)")
AddOption(cacheMutex, bool, true, "", 0, "Use a file lock to serialize access to the cache folder")
AddOption(ignoreCacheValid, bool, false, "", 0, "If set, allows to use RTC cached code files even if they are not valid for the current source code / parameters")
AddHelp("help", 'h')
EndConfig()

Expand Down
15 changes: 11 additions & 4 deletions 15 dependencies/FindO2GPU.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -111,13 +111,15 @@ if(ENABLE_CUDA)
endif()
endif()
if(CMAKE_CUDA_COMPILER)
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler \"${O2_GPU_CMAKE_CXX_FLAGS_NOSTD}\" --expt-relaxed-constexpr --extended-lambda --allow-unsupported-compiler -Xptxas -v -Xcompiler -Wno-attributes")
set(CMAKE_CUDA_FLAGS "-Xcompiler \"${O2_GPU_CMAKE_CXX_FLAGS_NOSTD}\" ${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --extended-lambda --allow-unsupported-compiler -Xptxas -v -Xcompiler -Wno-attributes")
if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "12.3")
string(APPEND CMAKE_CUDA_FLAGS " -Xcudafe --diag_suppress=20257") # TODO: Cleanup
endif()
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -lineinfo -Xcompiler \"${CMAKE_CXX_FLAGS_DEBUG}\" -Xptxas -O0 -Xcompiler -O0")
if(NOT CMAKE_BUILD_TYPE_UPPER STREQUAL "DEBUG")
set(CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -Xcompiler \"${CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER}}\" -Xptxas -O4 -Xcompiler -O4")
set(CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "-Xcompiler \"${CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER}}\" ${CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER}}")
if(CMAKE_BUILD_TYPE_UPPER STREQUAL "DEBUG")
set(CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -lineinfo -Xptxas -O0 -Xcompiler -O0")
else()
set(CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -Xptxas -O4 -Xcompiler -O4")
endif()
set(GPUCA_CUDA_NO_FAST_MATH_FLAGS "--ftz=false --prec-div=true --prec-sqrt=true --fmad false")
if(DEFINED GPUCA_NO_FAST_MATH AND "${GPUCA_NO_FAST_MATH}")
Expand Down Expand Up @@ -284,6 +286,11 @@ if(ENABLE_HIP)
endif()
set(CMAKE_HIP_FLAGS "${O2_GPU_CMAKE_CXX_FLAGS_NOSTD} ${CMAKE_HIP_FLAGS} ${O2_HIP_CMAKE_CXX_FLAGS}")
set(CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} ${CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER}}")
if(CMAKE_BUILD_TYPE_UPPER STREQUAL "DEBUG")
set(CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -O0 -ggdb")
else()
set(CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -O3")
endif()
else()
set(HIP_ENABLED OFF)
endif()
Expand Down
2 changes: 1 addition & 1 deletion 2 prodtests/full-system-test/dpl-workflow.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -320,7 +320,7 @@ if has_detector_calib PHS && workflow_has_parameter CALIB; then
fi

[[ ${O2_GPU_DOUBLE_PIPELINE:-$EPNSYNCMODE} == 1 ]] && GPU_CONFIG+=" --enableDoublePipeline"
[[ ${O2_GPU_RTC:-0} == 1 ]] && GPU_CONFIG_KEY+="GPU_proc_rtc.enable=1;GPU_proc_rtc.cacheOutput=1;GPU_proc.RTCcacheFolder=/tmp/o2_gpu_rtc_cache;"
[[ ${O2_GPU_RTC:-0} == 1 ]] && GPU_CONFIG_KEY+="GPU_proc_rtc.enable=1;GPU_proc_rtc.cacheOutput=1;GPU_proc.RTCcacheFolder=/var/tmp/o2_gpu_rtc_cache;"

( workflow_has_parameter AOD || [[ -z "$DISABLE_ROOT_OUTPUT" ]] || needs_root_output o2-emcal-cell-writer-workflow ) && has_detector EMC && RAW_EMC_SUBSPEC=" --subspecification 1 "
has_detector_reco MID && has_detector_matching MCHMID && MFTMCHConf="FwdMatching.useMIDMatch=true;" || MFTMCHConf="FwdMatching.useMIDMatch=false;"
Expand Down
Morty Proxy This is a proxified and sanitized view of the page, visit original site.