CMakeLists.txt 12.3 KB
Newer Older
Umang Yadav's avatar
Umang Yadav committed
1
# ####################################################################################
2
3
# The MIT License (MIT)
#
4
# Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
Umang Yadav's avatar
Umang Yadav committed
23
# ####################################################################################
Paul's avatar
Paul committed
24

25
find_package(hip REQUIRED)
26
if(NOT GPU_TARGETS)
27
28
29
30
31
    set(fatal_msg "HIP package is broken and has no GPU_TARGETS. Please pass GPU_TARGETS to cmake.")
    if(NOT WIN32)
        set(fatal_msg "${fatal_msg}\nUse -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*') to build for your GPU.")
    endif()
    message(FATAL_ERROR ${fatal_msg})
32
endif()
33
34
find_package(miopen REQUIRED)
message(STATUS "MIGraphX is using MIOpen")
Paul's avatar
Paul committed
35

36
# rocblas
37
38
find_package(rocblas REQUIRED)
message(STATUS "MIGraphX build with rocBLAS")
Paul's avatar
Paul committed
39

Artur Wojcik's avatar
Artur Wojcik committed
40
if(MIGRAPHX_USE_COMPOSABLEKERNEL)
41
42
    find_package(composable_kernel 1.0.0 REQUIRED COMPONENTS jit_library)
endif()
43

Chris Austen's avatar
Chris Austen committed
44
45
46
47
48
if(BUILD_DEV)
    set(MIGRAPHX_USE_HIPRTC OFF CACHE BOOL "Use hipRTC APIs")
else()
    set(MIGRAPHX_USE_HIPRTC ON CACHE BOOL "Use hipRTC APIs")
endif()
Umang Yadav's avatar
Umang Yadav committed
49

50
file(GLOB KERNEL_FILES CONFIGURE_DEPENDS
51
    ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/*.hpp)
52

Artur Wojcik's avatar
Artur Wojcik committed
53
if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
54
55
    list(REMOVE_ITEM KERNEL_FILES
        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck_gemm.hpp
Artur Wojcik's avatar
Artur Wojcik committed
56
        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck_gemm_softmax_gemm.hpp
57
58
59
60
        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck.hpp)
endif()

include(Embed)
61
add_embed_library(migraphx_kernels ${KERNEL_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/)
62

63
configure_file(device/targets.hpp.in include/migraphx/gpu/device/targets.hpp)
64
file(GLOB DEVICE_GPU_SRCS CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/device/*.cpp)
65
66
add_library(migraphx_device ${DEVICE_GPU_SRCS})

67
add_library(compile_for_gpu INTERFACE)
68
69
70
71
target_compile_features(compile_for_gpu INTERFACE cxx_std_17)
target_compile_options(compile_for_gpu INTERFACE -fno-gpu-rdc -Wno-cuda-compat -Wno-unused-command-line-argument -Xclang -fallow-half-arguments-and-returns)
target_link_options(compile_for_gpu INTERFACE  -fno-gpu-rdc -Wno-invalid-command-line-argument -Wno-unused-command-line-argument -Wno-option-ignored)
target_link_libraries(compile_for_gpu INTERFACE hip::device)
72
check_cxx_compiler_flag("--cuda-host-only -fhip-lambda-host-device -x hip" HAS_HIP_LAMBDA_HOST_DEVICE)
Umang Yadav's avatar
Umang Yadav committed
73

74
if(HAS_HIP_LAMBDA_HOST_DEVICE)
Umang Yadav's avatar
Umang Yadav committed
75
76
    message(STATUS "Enable -fhip-lambda-host-device")
    target_compile_options(compile_for_gpu INTERFACE -fhip-lambda-host-device)
77
endif()
78
79
80
81
82
83

set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
rocm_set_soversion(migraphx_device ${MIGRAPHX_SO_VERSION})
rocm_clang_tidy_check(migraphx_device)
target_link_libraries(migraphx_device PUBLIC migraphx)
target_link_libraries(migraphx_device PRIVATE compile_for_gpu)
Paul's avatar
Paul committed
84
target_include_directories(migraphx_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
85
target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_BINAR_DIR}/include>)
Paul's avatar
Paul committed
86
target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device/include>)
87
88
target_compile_options(migraphx_device PRIVATE -Wno-ignored-attributes)
migraphx_generate_export_header(migraphx_device DIRECTORY migraphx/gpu/device)
89

90
add_library(kernel_file_check EXCLUDE_FROM_ALL)
Umang Yadav's avatar
Umang Yadav committed
91

92
93
94
95
96
foreach(KERNEL_FILE ${KERNEL_FILES})
    get_filename_component(KERNEL_BASE_FILE ${KERNEL_FILE} NAME_WE)
    file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/kernels/include/migraphx/kernels/${KERNEL_BASE_FILE}.cpp "#include <migraphx/kernels/${KERNEL_BASE_FILE}.hpp>\n")
    target_sources(kernel_file_check PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/kernels/include/migraphx/kernels/${KERNEL_BASE_FILE}.cpp)
endforeach()
Umang Yadav's avatar
Umang Yadav committed
97

98
target_compile_definitions(kernel_file_check PRIVATE -DMIGRAPHX_NLOCAL=256)
99
100
101
102
103
target_include_directories(kernel_file_check PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/>)
target_link_libraries(kernel_file_check compile_for_gpu)

rocm_clang_tidy_check(kernel_file_check)

104
file(GLOB JIT_GPU_SRCS CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/jit/*.cpp)
105

Artur Wojcik's avatar
Artur Wojcik committed
106
107
108
109
if(NOT MIGRAPHX_USE_COMPOSABLEKERNEL)
    list(REMOVE_ITEM JIT_GPU_SRCS
            ${CMAKE_CURRENT_SOURCE_DIR}/jit/ck_gemm.cpp
            ${CMAKE_CURRENT_SOURCE_DIR}/jit/ck_gemm_softmax_gemm.cpp)
110
111
endif()

Paul's avatar
Paul committed
112
add_library(migraphx_gpu
Shucai Xiao's avatar
Shucai Xiao committed
113
    abs.cpp
114
    analyze_streams.cpp
115
    allocation_model.cpp
Shucai Xiao's avatar
Shucai Xiao committed
116
117
    argmax.cpp
    argmin.cpp
118
    code_object_op.cpp
119
    compile_ops.cpp
Paul Fultz II's avatar
Paul Fultz II committed
120
    compile_gen.cpp
Shucai Xiao's avatar
Shucai Xiao committed
121
    compile_hip.cpp
122
    compile_hip_code_object.cpp
123
    compile_miopen.cpp
124
    compiler.cpp
125
    device_name.cpp
126
    fuse_ck.cpp
Paul Fultz II's avatar
Paul Fultz II committed
127
    fuse_mlir.cpp
Paul's avatar
Paul committed
128
    fuse_ops.cpp
Shucai Xiao's avatar
Shucai Xiao committed
129
    gemm_impl.cpp
Paul's avatar
Paul committed
130
    hip.cpp
Shucai Xiao's avatar
Shucai Xiao committed
131
    kernel.cpp
Paul's avatar
Paul committed
132
    lowering.cpp
133
    logsoftmax.cpp
Shucai Xiao's avatar
Shucai Xiao committed
134
    loop.cpp
Khalique's avatar
Khalique committed
135
    lrn.cpp
Paul Fultz II's avatar
Paul Fultz II committed
136
    mlir.cpp
turneram's avatar
turneram committed
137
    multinomial.cpp
138
    no_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
139
    nonzero.cpp
Shucai Xiao's avatar
Shucai Xiao committed
140
    pack_args.cpp
141
    prefuse_ops.cpp
142
    perfdb.cpp
Shucai Xiao's avatar
Shucai Xiao committed
143
    pooling.cpp
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
144
    reverse.cpp
Shucai Xiao's avatar
Shucai Xiao committed
145
    rnn_variable_seq_lens.cpp
Shucai Xiao's avatar
Shucai Xiao committed
146
    rocblas.cpp
147
    scatter.cpp
Shucai Xiao's avatar
Shucai Xiao committed
148
    schedule_model.cpp
149
    sync_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
150
    target.cpp
151
    time_op.cpp
Shucai Xiao's avatar
Shucai Xiao committed
152
    topk.cpp
Shucai Xiao's avatar
Shucai Xiao committed
153
    write_literals.cpp
154
    ${JIT_GPU_SRCS}
Paul's avatar
Paul committed
155
)
156

Paul's avatar
Paul committed
157
set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
158
migraphx_generate_export_header(migraphx_gpu)
159

160
161
162
163
164
function(register_migraphx_gpu_ops PREFIX)
    foreach(OP ${ARGN})
        register_op(migraphx_gpu HEADER migraphx/gpu/${OP}.hpp OPERATORS gpu::${PREFIX}${OP} INCLUDES migraphx/gpu/context.hpp)
    endforeach()
endfunction()
Umang Yadav's avatar
Umang Yadav committed
165

166
167
168
169
register_migraphx_gpu_ops(hip_
    argmax
    argmin
    logsoftmax
Shucai Xiao's avatar
Shucai Xiao committed
170
    loop
turneram's avatar
turneram committed
171
    multinomial
Shucai Xiao's avatar
Shucai Xiao committed
172
    nonzero
173
    prefix_scan_sum
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
174
    reverse
175
    scatter
Shucai Xiao's avatar
Shucai Xiao committed
176
    topk
177
178
179
180
181
182
183
)
register_migraphx_gpu_ops(miopen_
    abs
    contiguous
    lrn
    pooling
)
Umang Yadav's avatar
Umang Yadav committed
184
185
register_op(migraphx_gpu
    HEADER migraphx/gpu/rnn_variable_seq_lens.hpp
186
187
    OPERATORS gpu::hip_rnn_var_sl_shift_sequence gpu::hip_rnn_var_sl_shift_output gpu::hip_rnn_var_sl_last_output
    INCLUDES migraphx/gpu/context.hpp)
Umang Yadav's avatar
Umang Yadav committed
188
189
register_op(migraphx_gpu
    HEADER migraphx/gpu/gemm.hpp
190
191
    OPERATORS gpu::rocblas_gemm<op::dot> gpu::rocblas_gemm<op::quant_dot>
    INCLUDES migraphx/gpu/context.hpp)
Umang Yadav's avatar
Umang Yadav committed
192
register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
193
    OPERATORS gpu::miopen_convolution<op::convolution> gpu::miopen_convolution<op::convolution_backwards> gpu::miopen_convolution<op::quant_convolution>
194
    INCLUDES migraphx/gpu/context.hpp)
Paul's avatar
Paul committed
195
rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
Paul's avatar
Paul committed
196
rocm_clang_tidy_check(migraphx_gpu)
Paul Fultz II's avatar
Paul Fultz II committed
197

198
set(MIGRAPHX_ENABLE_MLIR ON CACHE BOOL "")
Umang Yadav's avatar
Umang Yadav committed
199

200
if(MIGRAPHX_ENABLE_MLIR)
201
202
203
    # Find package rocMLIR
    find_package(rocMLIR 1.0.0 CONFIG REQUIRED)
    message(STATUS "Build with rocMLIR::rockCompiler ${rocMLIR_VERSION}")
Paul Fultz II's avatar
Paul Fultz II committed
204
    target_compile_definitions(migraphx_gpu PRIVATE "-DMIGRAPHX_MLIR")
205
206
207
    # Make this private to avoid multiple inclusions of LLVM symbols.
    # TODO: Fix rocMLIR's library to hide LLVM internals.
    target_link_libraries(migraphx_gpu PRIVATE rocMLIR::rockCompiler)
208
209
endif()

210
if(MIGRAPHX_USE_HIPRTC)
211
    find_package(hiprtc REQUIRED)
Umang Yadav's avatar
Umang Yadav committed
212
213
    message(STATUS "MIGraphX is using hipRTC")
    target_compile_definitions(migraphx_gpu PRIVATE -DMIGRAPHX_USE_HIPRTC=1)
214
    target_link_libraries(migraphx_gpu PUBLIC hiprtc::hiprtc)
215
else()
Umang Yadav's avatar
Umang Yadav committed
216
    message(STATUS "MIGraphX is using HIP Clang")
217

Umang Yadav's avatar
Umang Yadav committed
218
219
220
221
222
    # Get flags needed to compile hip
    include(TargetFlags)
    target_flags(HIP_COMPILER_FLAGS hip::device)

    # Remove cuda arch flags
223
224
    string(REGEX REPLACE "--cuda-gpu-arch=[a-z0-9]+ ?" "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
    string(REGEX REPLACE "--offload-arch=[a-z0-9:+-]+ ?" "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
Umang Yadav's avatar
Umang Yadav committed
225
226
227

    # Skip library paths since hip will incorrectly treat it as a source file
    string(APPEND HIP_COMPILER_FLAGS " ")
228

229
230
231
    if(WIN32)
        string(REPLACE "\\" "/" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
    endif()
Umang Yadav's avatar
Umang Yadav committed
232
233
234
235
    foreach(_unused RANGE 2)
        string(REGEX REPLACE " /[^ ]+\\.(a|so) " " " HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
    endforeach()

236
    message(STATUS "Hip compiler flags: \"${HIP_COMPILER_FLAGS}\"")
Umang Yadav's avatar
Umang Yadav committed
237
    target_compile_definitions(migraphx_gpu PRIVATE
238
239
        -DMIGRAPHX_HIP_COMPILER="${CMAKE_CXX_COMPILER}"
        -DMIGRAPHX_HIP_COMPILER_FLAGS="${HIP_COMPILER_FLAGS}"
Umang Yadav's avatar
Umang Yadav committed
240
241
242
    )

    if(DEFINED CMAKE_CXX_COMPILER_LAUNCHER)
243
244
245
246
247
        if(WIN32)
            execute_process(COMMAND where ${CMAKE_CXX_COMPILER_LAUNCHER} OUTPUT_VARIABLE MIGRAPHX_HIP_COMPILER_LAUNCHER)
        else()
            execute_process(COMMAND which ${CMAKE_CXX_COMPILER_LAUNCHER} OUTPUT_VARIABLE MIGRAPHX_HIP_COMPILER_LAUNCHER)
        endif()
Umang Yadav's avatar
Umang Yadav committed
248
        string(STRIP "${MIGRAPHX_HIP_COMPILER_LAUNCHER}" MIGRAPHX_HIP_COMPILER_LAUNCHER)
249
        target_compile_definitions(migraphx_gpu PRIVATE -DMIGRAPHX_HIP_COMPILER_LAUNCHER="${MIGRAPHX_HIP_COMPILER_LAUNCHER}")
Umang Yadav's avatar
Umang Yadav committed
250
    endif()
251
endif()
Paul Fultz II's avatar
Paul Fultz II committed
252

253
# Check miopen find mode api
254

255
256
include(CheckLibraryExists)
get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
257
get_target_property(ROCBLAS_LOCATION roc::rocblas LOCATION)
258
check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
259
check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
260
261
# Beta API for automated GEMM tuning
check_library_exists(roc::rocblas "rocblas_gemm_ex_get_solutions" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
262
263
# rocblas FP8 API
check_library_exists(roc::rocblas "rocblas_gemm_strided_batched_ex3" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_FP8_BETA_API)
264

Umang Yadav's avatar
Umang Yadav committed
265
set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
266

Umang Yadav's avatar
Umang Yadav committed
267
if(MIGRAPHX_USE_FIND_2_API)
268
269
270
271
272
273
    check_library_exists(MIOpen "miopenSetFindOptionPreallocatedTensor" "${MIOPEN_LOCATION}" HAS_PREALLOCATION_API)
    if(HAS_PREALLOCATION_API)
        target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API -DMIGRAPHX_PREALLOCATE_MIOPEN_BUFFERS)
    else()
        target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API)
    endif()
274
275
276
277
    message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
else()
    message(STATUS "MIGraphx is using legacy Find API in MIOpen")
endif()
278

279
280
if(HAS_FIND_MODE_API)
    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
281
    message(STATUS "MIGraphx is using Find Mode API of MIOpen")
282
283
284
285
else()
    message(STATUS "MIOpen does not have find mode api")
endif()

286
287
288
289
290
291
292
if(HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_USE_ROCBLAS_TUNING_API -DROCBLAS_BETA_FEATURES_API -DROCBLAS_NO_DEPRECATED_WARNINGS)
    message(STATUS "MIGraphx is using Beta API of rocBLAS")
else()
    message(STATUS "rocBLAS does not have User Tuning Beta API")
endif()

293
294
295
296
297
298
299
if(HAS_ROCBLAS_FP8_BETA_API)
    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_USE_ROCBLAS_FP8_API -DROCBLAS_BETA_FEATURES_API -DROCBLAS_NO_DEPRECATED_WARNINGS)
    message(STATUS "MIGraphX is using Beta API of rocBLAS for FP8 computations")
else()
    message(STATUS "rocBLAS does not have Fp8 Beta API")
endif()

Paul's avatar
Paul committed
300
target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
301
target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
Artur Wojcik's avatar
Artur Wojcik committed
302
if(MIGRAPHX_USE_COMPOSABLEKERNEL)
303
    target_link_libraries(migraphx_gpu PRIVATE composable_kernel::jit_library)
304
    target_compile_definitions(migraphx_gpu PRIVATE MIGRAPHX_USE_COMPOSABLEKERNEL=1)
305
endif()
mei-ye's avatar
mei-ye committed
306

307
add_subdirectory(driver)
308
add_subdirectory(hiprtc)
309

Paul's avatar
Paul committed
310
rocm_install_targets(
Umang Yadav's avatar
Umang Yadav committed
311
312
    TARGETS migraphx_gpu migraphx_device compile_for_gpu
    INCLUDE
Paul's avatar
Paul committed
313
314
    ${CMAKE_CURRENT_SOURCE_DIR}/include
)