Skip to content

Commit ca92bc4

Browse files
mconcasmconcas
committed
Improve HIP from CUDA generation
1 parent a1583a7 commit ca92bc4

6 files changed

Lines changed: 83 additions & 53 deletions

File tree

‎GPU/GPUbenchmark/CMakeLists.txt‎

Lines changed: 2 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -10,47 +10,11 @@
1010
# or submit itself to any jurisdiction.
1111

1212
if(CUDA_ENABLED)
13-
o2_add_executable(gpu-memory-benchmark-cuda
14-
SOURCES benchmark.cxx
15-
cuda/Kernels.cu
16-
PUBLIC_LINK_LIBRARIES Boost::program_options
17-
ROOT::Tree
18-
TARGETVARNAME targetName)
13+
add_subdirectory(cuda)
1914
endif()
2015

2116
if(HIP_ENABLED)
22-
# Hipify-perl
23-
set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl")
24-
25-
set(HIP_KERNEL "Kernels.hip.cxx")
26-
set(CU_KERNEL ${CMAKE_CURRENT_SOURCE_DIR}/cuda/Kernels.cu)
27-
set(HIP_KERNEL_PATH "${CMAKE_CURRENT_SOURCE_DIR}/hip/${HIP_KERNEL}")
28-
29-
if(EXISTS ${HIPIFY_EXECUTABLE})
30-
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${CU_KERNEL})
31-
add_custom_command(
32-
OUTPUT ${HIP_KERNEL_PATH}
33-
COMMAND ${HIPIFY_EXECUTABLE} --quiet-warnings ${CU_KERNEL} | sed '1{/\#include \"hip\\/hip_runtime.h\"/d}' > ${HIP_KERNEL_PATH}
34-
DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/cuda/Kernels.cu
35-
)
36-
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
37-
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
38-
set(CMAKE_CXX_EXTENSIONS OFF)
39-
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${O2_HIP_CMAKE_CXX_FLAGS} -fgpu-rdc")
40-
41-
o2_add_executable(gpu-memory-benchmark-hip
42-
SOURCES benchmark.cxx
43-
hip/Kernels.hip.cxx
44-
PUBLIC_LINK_LIBRARIES hip::host
45-
Boost::program_options
46-
ROOT::Tree
47-
TARGETVARNAME targetName)
48-
49-
if(HIP_AMDGPUTARGET)
50-
# Need to add gpu target also to link flags due to gpu-rdc option
51-
target_link_options(${targetName} PUBLIC --amdgpu-target=${HIP_AMDGPUTARGET})
52-
endif()
53-
endif()
17+
add_subdirectory(hip)
5418
endif()
5519

5620
o2_add_test_root_macro(macro/showBenchmarks.C)

‎GPU/GPUbenchmark/Shared/Utils.h‎

Lines changed: 0 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -192,19 +192,6 @@ class BSDRnd : public LCGRnd
192192
__host__ __device__ int rnd() { return LCGRnd::rnd(); }
193193
};
194194

195-
196-
// CUDA does not support <type4> operations:
197-
// https://forums.developer.nvidia.com/t/swizzling-float4-arithmetic-support/217
198-
#ifndef __HIPCC__
199-
inline __host__ __device__ void operator+=(int4 &a, int4 b)
200-
{
201-
a.x += b.x;
202-
a.y += b.y;
203-
a.z += b.z;
204-
a.w += b.w;
205-
}
206-
#endif
207-
208195
namespace o2
209196
{
210197
namespace benchmark

‎GPU/GPUbenchmark/cuda/CMakeLists.txt‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
# Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2+
# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3+
# All rights not expressly granted are reserved.
4+
#
5+
# This software is distributed under the terms of the GNU General Public
6+
# License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7+
#
8+
# In applying this license CERN does not waive the privileges and immunities
9+
# granted to it by virtue of its status as an Intergovernmental Organization
10+
# or submit itself to any jurisdiction.
11+
12+
if(CUDA_ENABLED)
13+
message(STATUS "Building GPU CUDA benchmark")
14+
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
15+
o2_add_executable(gpu-memory-benchmark-cuda
16+
SOURCES benchmark.cu
17+
Kernels.cu
18+
PUBLIC_LINK_LIBRARIES Boost::program_options
19+
ROOT::Tree
20+
TARGETVARNAME targetName)
21+
set_target_properties(${targeName} PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
22+
endif()

‎GPU/GPUbenchmark/cuda/Kernels.cu‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,19 @@
3131
double bytesToconfig(size_t s) { return (double)s / (1024.0); }
3232
double bytesToGB(size_t s) { return (double)s / GB; }
3333

34+
35+
// CUDA does not support <type4> operations:
36+
// https://forums.developer.nvidia.com/t/swizzling-float4-arithmetic-support/217
37+
#ifndef __HIPCC__
38+
inline __host__ __device__ void operator+=(int4 &a, int4 b)
39+
{
40+
a.x += b.x;
41+
a.y += b.y;
42+
a.z += b.z;
43+
a.w += b.w;
44+
}
45+
#endif
46+
3447
namespace o2
3548
{
3649
namespace benchmark

GPU/GPUbenchmark/benchmark.cxx renamed to GPU/GPUbenchmark/cuda/benchmark.cu

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,10 +9,13 @@
99
// granted to it by virtue of its status as an Intergovernmental Organization
1010
// or submit itself to any jurisdiction.
1111
///
12-
/// \file benchmark.cxx
12+
/// \file benchmark.cu
1313
/// \author mconcas@cern.ch
1414
///
15-
#include "Shared/Kernels.h"
15+
#include "../Shared/Kernels.h"
16+
#if defined(__HIPCC__)
17+
#include "hip/hip_runtime.h"
18+
#endif
1619
#define VERSION "version 0.1-latest-#6773"
1720

1821
bool parseArgs(o2::benchmark::benchmarkOpts& conf, int argc, const char* argv[])

‎GPU/GPUbenchmark/hip/CMakeLists.txt‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
# Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2+
# See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3+
# All rights not expressly granted are reserved.
4+
#
5+
# This software is distributed under the terms of the GNU General Public
6+
# License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7+
#
8+
# In applying this license CERN does not waive the privileges and immunities
9+
# granted to it by virtue of its status as an Intergovernmental Organization
10+
# or submit itself to any jurisdiction.
11+
12+
message(STATUS "Building GPU HIP benchmark")
13+
# Hipify-perl to generate HIP sources
14+
set(HIPIFY_EXECUTABLE "/opt/rocm/bin/hipify-perl")
15+
file(GLOB CUDA_SOURCES_FULL_PATH "../cuda/*.cu")
16+
foreach(file ${CUDA_SOURCES_FULL_PATH})
17+
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${file})
18+
get_filename_component(CUDA_SOURCE ${file} NAME)
19+
string(REPLACE ".cu" "" CUDA_SOURCE_NAME ${CUDA_SOURCE})
20+
add_custom_command(
21+
OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/${CUDA_SOURCE_NAME}.hip.cxx
22+
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.cxx
23+
DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../cuda/${CUDA_SOURCE}
24+
)
25+
endforeach()
26+
27+
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
28+
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
29+
set(CMAKE_CXX_EXTENSIONS OFF)
30+
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${O2_HIP_CMAKE_CXX_FLAGS} -fgpu-rdc")
31+
o2_add_executable(gpu-memory-benchmark-hip
32+
SOURCES benchmark.hip.cxx
33+
Kernels.hip.cxx
34+
PUBLIC_LINK_LIBRARIES hip::host
35+
Boost::program_options
36+
ROOT::Tree
37+
TARGETVARNAME targetName)
38+
if(HIP_AMDGPUTARGET)
39+
# Need to add gpu target also to link flags due to gpu-rdc option
40+
target_link_options(${targetName} PUBLIC --amdgpu-target=${HIP_AMDGPUTARGET})
41+
endif()

0 commit comments

Comments
 (0)