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
130
    gather.cpp
    gemm_impl.cpp
Paul's avatar
Paul committed
131
    hip.cpp
Shucai Xiao's avatar
Shucai Xiao committed
132
    kernel.cpp
Paul's avatar
Paul committed
133
    lowering.cpp
134
    logsoftmax.cpp
Shucai Xiao's avatar
Shucai Xiao committed
135
    loop.cpp
Khalique's avatar
Khalique committed
136
    lrn.cpp
Paul Fultz II's avatar
Paul Fultz II committed
137
    mlir.cpp
turneram's avatar
turneram committed
138
    multinomial.cpp
139
    no_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
140
    nonzero.cpp
Shucai Xiao's avatar
Shucai Xiao committed
141
    pack_args.cpp
142
    prefuse_ops.cpp
Shucai Xiao's avatar
Shucai Xiao committed
143
    pad.cpp
144
    perfdb.cpp
Shucai Xiao's avatar
Shucai Xiao committed
145
    pooling.cpp
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
146
    reverse.cpp
Shucai Xiao's avatar
Shucai Xiao committed
147
    rnn_variable_seq_lens.cpp
Shucai Xiao's avatar
Shucai Xiao committed
148
    rocblas.cpp
149
    scatter.cpp
Shucai Xiao's avatar
Shucai Xiao committed
150
    schedule_model.cpp
151
    sync_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
152
    target.cpp
153
    time_op.cpp
Shucai Xiao's avatar
Shucai Xiao committed
154
    topk.cpp
Shucai Xiao's avatar
Shucai Xiao committed
155
    write_literals.cpp
156
    ${JIT_GPU_SRCS}
Paul's avatar
Paul committed
157
)
158

Paul's avatar
Paul committed
159
set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
160
migraphx_generate_export_header(migraphx_gpu)
161

162
163
164
165
166
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
167

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

202
set(MIGRAPHX_ENABLE_MLIR ON CACHE BOOL "")
Umang Yadav's avatar
Umang Yadav committed
203

204
if(MIGRAPHX_ENABLE_MLIR)
205
206
207
    # 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
208
    target_compile_definitions(migraphx_gpu PRIVATE "-DMIGRAPHX_MLIR")
209
210
211
    # 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)
212
213
endif()

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

Umang Yadav's avatar
Umang Yadav committed
222
223
224
225
226
    # Get flags needed to compile hip
    include(TargetFlags)
    target_flags(HIP_COMPILER_FLAGS hip::device)

    # Remove cuda arch flags
227
228
    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
229
230
231

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

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

240
    message(STATUS "Hip compiler flags: \"${HIP_COMPILER_FLAGS}\"")
Umang Yadav's avatar
Umang Yadav committed
241
    target_compile_definitions(migraphx_gpu PRIVATE
242
243
        -DMIGRAPHX_HIP_COMPILER="${CMAKE_CXX_COMPILER}"
        -DMIGRAPHX_HIP_COMPILER_FLAGS="${HIP_COMPILER_FLAGS}"
Umang Yadav's avatar
Umang Yadav committed
244
245
246
    )

    if(DEFINED CMAKE_CXX_COMPILER_LAUNCHER)
247
248
249
250
251
        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
252
        string(STRIP "${MIGRAPHX_HIP_COMPILER_LAUNCHER}" MIGRAPHX_HIP_COMPILER_LAUNCHER)
253
        target_compile_definitions(migraphx_gpu PRIVATE -DMIGRAPHX_HIP_COMPILER_LAUNCHER="${MIGRAPHX_HIP_COMPILER_LAUNCHER}")
Umang Yadav's avatar
Umang Yadav committed
254
    endif()
255
endif()
Paul Fultz II's avatar
Paul Fultz II committed
256

257
# Check miopen find mode api
258

259
260
include(CheckLibraryExists)
get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
261
get_target_property(ROCBLAS_LOCATION roc::rocblas LOCATION)
262
check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
263
check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
264
265
# Beta API for automated GEMM tuning
check_library_exists(roc::rocblas "rocblas_gemm_ex_get_solutions" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_TUNING_BETA_FEATURE_API)
Umang Yadav's avatar
Umang Yadav committed
266
# rocblas FP8 API
Umang Yadav's avatar
Umang Yadav committed
267
check_library_exists(roc::rocblas "rocblas_gemm_strided_batched_ex3" "${ROCBLAS_LOCATION}" HAS_ROCBLAS_FP8_BETA_API)
268

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

Umang Yadav's avatar
Umang Yadav committed
271
if(MIGRAPHX_USE_FIND_2_API)
272
273
274
275
276
277
    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()
278
279
280
281
    message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
else()
    message(STATUS "MIGraphx is using legacy Find API in MIOpen")
endif()
282

283
284
if(HAS_FIND_MODE_API)
    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
285
    message(STATUS "MIGraphx is using Find Mode API of MIOpen")
286
287
288
289
else()
    message(STATUS "MIOpen does not have find mode api")
endif()

290
291
292
293
294
295
296
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()

Umang Yadav's avatar
Umang Yadav committed
297
298
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)
Umang Yadav's avatar
Umang Yadav committed
299
    message(STATUS "MIGraphX is using Beta API of rocBLAS for FP8 computations")
Umang Yadav's avatar
Umang Yadav committed
300
else()
Umang Yadav's avatar
Umang Yadav committed
301
    message(STATUS "rocBLAS does not have Fp8 Beta API")
Umang Yadav's avatar
Umang Yadav committed
302
303
endif()

Paul's avatar
Paul committed
304
target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
305
target_link_libraries(migraphx_gpu PRIVATE migraphx_device migraphx_kernels)
Artur Wojcik's avatar
Artur Wojcik committed
306
if(MIGRAPHX_USE_COMPOSABLEKERNEL)
307
308
    target_link_libraries(migraphx_gpu PRIVATE composable_kernel::jit_library)
endif()
mei-ye's avatar
mei-ye committed
309

310
add_subdirectory(driver)
311
add_subdirectory(hiprtc)
312

Paul's avatar
Paul committed
313
rocm_install_targets(
Umang Yadav's avatar
Umang Yadav committed
314
315
    TARGETS migraphx_gpu migraphx_device compile_for_gpu
    INCLUDE
Paul's avatar
Paul committed
316
317
    ${CMAKE_CURRENT_SOURCE_DIR}/include
)