Skip to content

Commit 3102627

Browse files
davidrohrdavidrohr
committed
GPU: Reorganize some files, split OCL code in kernel and non-kernel related parts
1 parent 9c90527 commit 3102627

16 files changed

+351
-298
lines changed

‎GPU/GPUTracking/Base/cuda/CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@ endif()
1818
message(STATUS "Building GPUTracking with CUDA support ${TMP_TARGET}")
1919

2020
set(SRCS GPUReconstructionCUDA.cu GPUReconstructionCUDAGenRTC.cxx GPUReconstructionCUDAKernels.cu)
21-
set(HDRS GPUReconstructionCUDA.h GPUReconstructionCUDAInternals.h GPUReconstructionCUDADef.h GPUReconstructionCUDAIncludes.h CUDAThrustHelpers.h)
21+
set(HDRS GPUReconstructionCUDA.h GPUReconstructionCUDAInternals.h GPUReconstructionCUDADef.h GPUReconstructionCUDAIncludesHost.h CUDAThrustHelpers.h)
2222
# -------------------------------- Prepare RTC -------------------------------------------------------
2323
enable_language(ASM)
2424
if(ALIGPU_BUILD_TYPE STREQUAL "O2")
@@ -67,7 +67,7 @@ set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionCUDArtc)
6767
# cmake-format: off
6868
add_custom_command(
6969
OUTPUT ${GPU_RTC_BIN}.src
70-
COMMAND cat ${GPUDIR}/Base/cuda/GPUReconstructionCUDAIncludes.h > ${GPU_RTC_BIN}.src
70+
COMMAND cp ${GPUDIR}/Base/cuda/GPUReconstructionCUDAIncludesHost.h ${GPU_RTC_BIN}.src
7171
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
7272
MAIN_DEPENDENCY ${GPU_RTC_SRC}
7373
IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC}

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,7 @@
1313
/// \author David Rohr
1414

1515
#define GPUCA_GPUCODE_HOSTONLY
16-
#include "GPUReconstructionCUDADef.h"
17-
#include "GPUReconstructionCUDAIncludes.h"
16+
#include "GPUReconstructionCUDAIncludesHost.h"
1817

1918
#include <cuda_profiler_api.h>
2019

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDADef.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
// granted to it by virtue of its status as an Intergovernmental Organization
1010
// or submit itself to any jurisdiction.
1111

12-
/// \file GPUReconstructionCUDDef.h
12+
/// \file GPUReconstructionCUDADef.h
1313
/// \author David Rohr
1414

1515
#ifndef O2_GPU_GPURECONSTRUCTIONCUDADEF_H

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAExternalProvider.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,7 @@
1212
/// \file GPUReconstructionCUDAExternalProvider.cu
1313
/// \author David Rohr
1414

15-
#include "GPUReconstructionCUDADef.h"
16-
#include "GPUReconstructionCUDAIncludes.h"
15+
#include "GPUReconstructionCUDAIncludesHost.h"
1716

1817
#include "GPUReconstructionCUDA.h"
1918
#include "GPUReconstructionCUDAInternals.h"

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAGenRTC.cxx‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,9 @@ QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command_arch);
3232

3333
int32_t GPUReconstructionCUDA::genRTC(std::string& filename, uint32_t& nCompile)
3434
{
35-
std::string rtcparam = std::string(mProcessingSettings.rtc.optSpecialCode ? "#define GPUCA_RTC_SPECIAL_CODE(...) __VA_ARGS__\n" : "#define GPUCA_RTC_SPECIAL_CODE(...)\n") + GPUParamRTC::generateRTCCode(param(), mProcessingSettings.rtc.optConstexpr);
35+
std::string rtcparam = std::string("#define GPUCA_RTC_CODE\n") +
36+
std::string(mProcessingSettings.rtc.optSpecialCode ? "#define GPUCA_RTC_SPECIAL_CODE(...) __VA_ARGS__\n" : "#define GPUCA_RTC_SPECIAL_CODE(...)\n") +
37+
GPUParamRTC::generateRTCCode(param(), mProcessingSettings.rtc.optConstexpr);
3638
if (filename == "") {
3739
filename = "/tmp/o2cagpu_rtc_";
3840
}

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAIncludes.h renamed to GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAIncludesHost.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,4 +32,8 @@
3232
#include <sm_20_atomic_functions.h>
3333
#include <cuda_fp16.h>
3434

35+
#ifndef GPUCA_RTC_CODE
36+
#include "GPUReconstructionCUDADef.h"
37+
#endif
38+
3539
#endif

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernels.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,7 @@
1212
/// \file GPUReconstructionCUDAKernels.cu
1313
/// \author David Rohr
1414

15-
#include "GPUReconstructionCUDADef.h"
16-
#include "GPUReconstructionCUDAIncludes.h"
15+
#include "GPUReconstructionCUDAIncludesHost.h"
1716

1817
#include "GPUReconstructionCUDA.h"
1918
#include "GPUReconstructionCUDAInternals.h"

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAkernel.template.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,7 @@
1313
/// \author David Rohr
1414

1515
#define GPUCA_GPUCODE_COMPILEKERNELS
16-
#include "GPUReconstructionCUDAIncludes.h"
17-
#include "GPUReconstructionCUDADef.h"
16+
#include "GPUReconstructionCUDAIncludesHost.h"
1817
#define GPUCA_KRNL_REG(args) __launch_bounds__(GPUCA_M_MAX2_3(GPUCA_M_STRIP(args)))
1918
#define GPUCA_KRNL(...) GPUCA_KRNL_WRAP(GPUCA_KRNL_LOAD_, __VA_ARGS__)
2019
#define GPUCA_KRNL_LOAD_single(...) GPUCA_KRNLGPU_SINGLE(__VA_ARGS__);

‎GPU/GPUTracking/Base/hip/CMakeLists.txt‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ if(NOT DEFINED GPUCA_HIP_HIPIFY_FROM_CUDA OR "${GPUCA_HIP_HIPIFY_FROM_CUDA}")
2525
set(GPUCA_HIP_SOURCE_DIR ${CMAKE_CURRENT_BINARY_DIR}/hipify)
2626
file(MAKE_DIRECTORY ${GPUCA_HIP_SOURCE_DIR})
2727
set(GPUCA_HIP_FILE_LIST GPUReconstructionCUDA.cu GPUReconstructionCUDAExternalProvider.cu GPUReconstructionCUDA.h GPUReconstructionCUDAInternals.h GPUReconstructionCUDAkernel.template.cu CUDAThrustHelpers.h GPUReconstructionCUDADef.h GPUReconstructionCUDAGenRTC.cxx GPUReconstructionCUDAKernels.cu GPUReconstructionCUDArtc.cu)
28-
set(GPUCA_HIP_LOCAL_FILE_LIST GPUReconstructionHIPIncludes.h)
28+
set(GPUCA_HIP_LOCAL_FILE_LIST GPUReconstructionHIPIncludesHost.h)
2929
set(HIP_SOURCES "")
3030
foreach(file ${GPUCA_HIP_FILE_LIST})
3131
get_filename_component(ABS_CUDA_SORUCE ../cuda/${file} ABSOLUTE)
@@ -63,7 +63,7 @@ endif()
6363

6464
set(SRCS ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIP.hip ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPKernels.hip)
6565
set(SRCS_CXX ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPGenRTC.cxx)
66-
set(HDRS ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIP.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPInternals.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPDef.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPIncludes.h ${GPUCA_HIP_SOURCE_DIR}/HIPThrustHelpers.h)
66+
set(HDRS ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIP.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPInternals.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPDef.h ${GPUCA_HIP_SOURCE_DIR}/GPUReconstructionHIPIncludesHost.h ${GPUCA_HIP_SOURCE_DIR}/HIPThrustHelpers.h)
6767

6868
# -------------------------------- Prepare RTC -------------------------------------------------------
6969
enable_language(ASM)
@@ -104,7 +104,7 @@ set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionHIPrtc)
104104
# cmake-format: off
105105
add_custom_command(
106106
OUTPUT ${GPU_RTC_BIN}.src
107-
COMMAND cat ${GPUDIR}/Base/hip/GPUReconstructionHIPIncludes.h > ${GPU_RTC_BIN}.src
107+
COMMAND cp ${GPUDIR}/Base/hip/GPUReconstructionHIPIncludesHost.h ${GPU_RTC_BIN}.src
108108
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
109109
MAIN_DEPENDENCY ${GPU_RTC_SRC}
110110
IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC}

GPU/GPUTracking/Base/hip/GPUReconstructionHIPIncludes.h renamed to GPU/GPUTracking/Base/hip/GPUReconstructionHIPIncludesHost.h

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
// granted to it by virtue of its status as an Intergovernmental Organization
1010
// or submit itself to any jurisdiction.
1111

12-
/// \file GPUReconstructionHIPInclude.h
12+
/// \file GPUReconstructionHIPIncludesHost.h
1313
/// \author David Rohr
1414

1515
#ifndef O2_GPU_RECONSTRUCTIONHIPINCLUDES_H
@@ -27,4 +27,8 @@
2727
#include <thrust/device_ptr.h>
2828
#pragma GCC diagnostic pop
2929

30+
#ifndef GPUCA_RTC_CODE
31+
#include "GPUReconstructionHIPDef.h"
32+
#endif
33+
3034
#endif

0 commit comments

Comments
 (0)