GPU: Some work to prepare using <type_traits> and <array> headers from system for GPU

Author: davidrohr

Common/MathUtils/include/MathUtils/SMatrixGPU.h

@@ -29,7 +29,9 @@
 #include "GPUCommonMath.h"
 #include "GPUCommonAlgorithm.h"
 #include "GPUCommonLogger.h"
-#include "GPUCommonTypeTraits.h"
+#ifndef GPUCA_GPUCODE_DEVICE
+#include <type_traits>
+#endif
 
 namespace o2::math_utils::detail
 {

DataFormats/common/include/CommonDataFormat/AbstractRef.h

@@ -18,7 +18,10 @@
 
 #include "GPUCommonDef.h"
 #include "GPUCommonRtypes.h"
-#include "GPUCommonTypeTraits.h"
+#ifndef GPUCA_GPUCODE_DEVICE
+#include <type_traits>
+#endif
+
 
 namespace o2::dataformats
 {

Detectors/Raw/include/DetectorsRaw/RDHUtils.h

@@ -19,13 +19,15 @@
 #include "GPUCommonRtypes.h"
 #include "Headers/RAWDataHeader.h"
 #include "Headers/RDHAny.h"
-#include "GPUCommonTypeTraits.h"
+#ifndef GPUCA_GPUCODE_DEVICE
+#include <type_traits>
+#endif
 #if !defined(GPUCA_GPUCODE)
 #include "CommonDataFormat/InteractionRecord.h"
 #endif
 #if !defined(GPUCA_GPUCODE) && !defined(GPUCA_STANDALONE)
 #include "Headers/DAQID.h"
-#endif // GPUCA_GPUCODE / GPUCA_STANDALONE
+#endif
 
 namespace o2
 {

GPU/GPUTracking/Base/GPUStdSystemHeaders.h

@@ -0,0 +1,25 @@
+// 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.
+
+/// \file GPUStdSystemHeaders.h
+/// \author David Rohr
+
+#ifndef GPUSTDSYSTEMHEADERS_H
+#define GPUSTDSYSTEMHEADERS_H
+
+#include <string>
+#include <cstddef>
+#include <cstdint>
+#include <cfloat>
+#include <cmath>
+#include <type_traits>
+
+#endif

GPU/GPUTracking/Base/cuda/CMakeLists.txt

@@ -67,7 +67,9 @@ set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionCUDArtc)
 # cmake-format: off
 add_custom_command(
     OUTPUT ${GPU_RTC_BIN}.src
-    COMMAND cp ${GPUDIR}/Base/cuda/GPUReconstructionCUDAIncludesSystem.h ${GPU_RTC_BIN}.src
+    COMMAND cp ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPU_RTC_BIN}.src
+    COMMAND cat ${GPUDIR}/Base/cuda/GPUReconstructionCUDAIncludesSystem.h | grep -v GPUStdSystemHeaders.h >> ${GPU_RTC_BIN}.src
+    COMMAND cat ${GPUDIR}/Base/GPUStdSystemHeaders.h >> ${GPU_RTC_BIN}.src
     COMMAND ${CMAKE_CUDA_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_CUDA_STANDARD} -D__CUDA_ARCH__=${RTC_CUDA_ARCH} -D__CUDACC__ -x c++ -E -Xcompiler "-nostdinc -P" ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src
     MAIN_DEPENDENCY ${GPU_RTC_SRC}
     IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC}

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAIncludesSystem.h

@@ -15,9 +15,10 @@
 #ifndef O2_GPU_GPURECONSTRUCTIONCUDAINCLUDESSYSTEM_H
 #define O2_GPU_GPURECONSTRUCTIONCUDAINCLUDESSYSTEM_H
 
-#include <cstdint>
-#include <type_traits>
-#include <string>
+#ifndef GPUCA_GPUCODE_GENRTC
+#include "GPUStdSystemHeaders.h"
+#endif
+
 #include <cuda_runtime.h>
 #include <cuda.h>
 #include <cooperative_groups.h>

GPU/GPUTracking/Base/hip/CMakeLists.txt

@@ -114,7 +114,8 @@ set(GPU_RTC_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionHIPrtc)
 # cmake-format: off
 add_custom_command(
     OUTPUT ${GPU_RTC_BIN}.src
-    COMMAND cp ${GPUDIR}/Base/hip/GPUReconstructionHIPIncludesSystem.h ${GPU_RTC_BIN}.src
+    COMMAND cp ${GPUDIR}/Base/GPUStdSystemHeaders.h ${GPU_RTC_BIN}.src
+    COMMAND cat ${GPUDIR}/Base/hip/GPUReconstructionHIPIncludesSystem.h | grep -v GPUStdSystemHeaders.h >> ${GPU_RTC_BIN}.src
     COMMAND ${CMAKE_HIP_COMPILER} ${GPU_RTC_DEFINES} ${GPU_RTC_INCLUDES} -std=c++${CMAKE_HIP_STANDARD} -D__HIPCC__ -D__HIP_DEVICE_COMPILE__ -x c++ -nostdinc -E -P ${GPU_RTC_SRC} >> ${GPU_RTC_BIN}.src
     MAIN_DEPENDENCY ${GPU_RTC_SRC}
     IMPLICIT_DEPENDS CXX ${GPU_RTC_SRC}

GPU/GPUTracking/Base/hip/GPUReconstructionHIPIncludesSystem.h

@@ -15,6 +15,10 @@
 #ifndef O2_GPU_RECONSTRUCTIONHIPINCLUDESSYSTEM_H
 #define O2_GPU_RECONSTRUCTIONHIPINCLUDESSYSTEM_H
 
+#ifndef GPUCA_GPUCODE_GENRTC
+#include "GPUStdSystemHeaders.h"
+#endif
+
 #include <hip/hip_runtime.h>
 #include <hip/hip_ext.h>
 #include <hipcub/hipcub.hpp>

GPU/GPUTracking/Base/opencl/CMakeLists.txt

@@ -23,7 +23,7 @@ endif()
 set(CL_SRC ${GPUDIR}/Base/opencl/GPUReconstructionOCL.cl)
 set(CL_BIN ${CMAKE_CURRENT_BINARY_DIR}/GPUReconstructionOCLCode)
 
-set(OCL_FLAGS -Dcl_clang_storage_class_specifiers -cl-std=CLC++2021 ${GPUCA_OCL_DENORMALS_FLAGS})
+set(OCL_FLAGS -Dcl_clang_storage_class_specifiers -x cl -cl-std=CLC++2021 ${GPUCA_OCL_DENORMALS_FLAGS})
 if(NOT GPUCA_DETERMINISTIC_MODE GREATER_EQUAL ${GPUCA_DETERMINISTIC_MODE_MAP_NO_FAST_MATH})
   set(OCL_FLAGS ${OCL_FLAGS} -cl-mad-enable -cl-no-signed-zeros -cl-fast-relaxed-math)
 else()
@@ -43,6 +43,11 @@ if (NOT DEFINED GPUCA_OCL_SPIRV_VERSION)
   set(GPUCA_OCL_SPIRV_VERSION 1.2)
 endif()
 
+# execute_process(COMMAND bash -c "${LLVM_CLANG} -stdlib=libc++ -E -H -x c++ - <<< '#include <type_traits>' 2>&1 1>/dev/null | grep type_traits | head -n 1 | sed 's/^\\.* *//'"
+#                 OUTPUT_VARIABLE CLANG_STD_INCLUDE_DIR)
+# get_filename_component(CLANG_STD_INCLUDE_DIR "${CLANG_STD_INCLUDE_DIR}" DIRECTORY)
+# get_filename_component(CLANG_STD_INCLUDE_DIR "${CLANG_STD_INCLUDE_DIR}" ABSOLUTE) # TODO: For using <type_traits> in OpenCL, we would need to add -I${CLANG_STD_INCLUDE_DIR}
+
 if(OPENCL_ENABLED_SPIRV) # BUILD OpenCL intermediate code for SPIR-V target
   # executes clang to create llvm IL code
   # Add -fintegrated-objemitter once we switch to clang >= 17
@@ -71,6 +76,7 @@ if(OPENCL_ENABLED) # BUILD OpenCL source code for runtime compilation target
   add_custom_command(
       OUTPUT ${CL_BIN}.src
       COMMAND ${LLVM_CLANG}
+              -target spir64
               -Wno-unused-command-line-argument
               ${OCL_FLAGS}
               ${OCL_DEFINECL}

GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cl

@@ -70,6 +70,8 @@ typedef signed char int8_t;
 #endif
 #define assert(param)
 
+#include "GPUCommonDef.h"
+#include "GPUCommonTypeTraits.h" // TODO: Once possible in OpenCL, should use GPUStdSystemHeaders.h here
 #include "GPUConstantMem.h"
 #include "GPUReconstructionIncludesDeviceAll.h"