CMakeLists.txt 11.6 KB
Newer Older
Umang Yadav's avatar
Umang Yadav committed
1
# ####################################################################################
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
# The MIT License (MIT)
#
# Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
#
# 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

Artur Wojcik's avatar
cmakes  
Artur Wojcik committed
25
list(APPEND CMAKE_PREFIX_PATH /opt/rocm $ENV{ROCM_PATH})
Paul's avatar
Paul committed
26

Artur Wojcik's avatar
cmakes  
Artur Wojcik committed
27
28
find_package(miopen REQUIRED)
message(STATUS "MIGraphX is using MIOpen")
29

Artur Wojcik's avatar
cmakes  
Artur Wojcik committed
30
31
32
# rocblas
find_package(rocblas REQUIRED)
message(STATUS "MIGraphX build with rocBLAS")
Paul's avatar
Paul committed
33

34
35
36
37
if(NOT WIN32)
    # TODO: re-enable when CK is ported to Windows
    find_package(composable_kernel 1.0.0 REQUIRED COMPONENTS jit_library)
endif()
38

Chris Austen's avatar
Chris Austen committed
39
40
41
42
43
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
44

45
file(GLOB KERNEL_FILES CONFIGURE_DEPENDS
46
    ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/*.hpp)
47
48
49
50
51
52
53
54
55
56
57
58

if(WIN32)
    # TODO: re-enable when CK is ported to Windows
    list(REMOVE_ITEM KERNEL_FILES
        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck_gemm.hpp
        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck.hpp)
endif()

if(NOT WIN32)
    include(Embed)
    add_embed_library(migraphx_kernels ${KERNEL_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/)
endif()
59

60
configure_file(device/targets.hpp.in include/migraphx/gpu/device/targets.hpp)
61
add_library(compile_for_gpu INTERFACE)
Artur Wojcik's avatar
cmakes  
Artur Wojcik committed
62
63
64
65
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)
66
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
67

68
if(HAS_HIP_LAMBDA_HOST_DEVICE)
Umang Yadav's avatar
Umang Yadav committed
69
70
    message(STATUS "Enable -fhip-lambda-host-device")
    target_compile_options(compile_for_gpu INTERFACE -fhip-lambda-host-device)
71
endif()
72

Artur Wojcik's avatar
cmakes  
Artur Wojcik committed
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
add_library(migraphx_device
        device/argmax.cpp
        device/argmin.cpp
        device/contiguous.cpp
        device/fill.cpp
        device/gather.cpp
        device/int8_gemm_pack.cpp
        device/logsoftmax.cpp
        device/multinomial.cpp
        device/nonzero.cpp
        device/pad.cpp
        device/prefix_scan_sum.cpp
        device/reverse.cpp
        device/rnn_variable_seq_lens.cpp
        device/scatter.cpp
        device/topk.cpp)

90
91
92
93
94
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
95
target_include_directories(migraphx_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
96
target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_BINAR_DIR}/include>)
Paul's avatar
Paul committed
97
target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device/include>)
98
99
target_compile_options(migraphx_device PRIVATE -Wno-ignored-attributes)
migraphx_generate_export_header(migraphx_device DIRECTORY migraphx/gpu/device)
100

101
add_library(kernel_file_check EXCLUDE_FROM_ALL)
Umang Yadav's avatar
Umang Yadav committed
102

103
104
105
106
107
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
108

109
target_compile_definitions(kernel_file_check PRIVATE -DMIGRAPHX_NLOCAL=256)
110
111
112
113
114
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)

115
file(GLOB JIT_GPU_SRCS CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/jit/*.cpp)
116

turneram's avatar
turneram committed
117
if(WIN32)
118
119
120
121
    # TODO: re-enable when CK is ported to Windows
    list(REMOVE_ITEM JIT_GPU_SRCS ${CMAKE_CURRENT_SOURCE_DIR}/jit/ck_gemm.cpp)
endif()

Paul's avatar
Paul committed
122
add_library(migraphx_gpu
Shucai Xiao's avatar
Shucai Xiao committed
123
    abs.cpp
124
    analyze_streams.cpp
125
    allocation_model.cpp
Shucai Xiao's avatar
Shucai Xiao committed
126
127
    argmax.cpp
    argmin.cpp
128
    code_object_op.cpp
129
    compile_ops.cpp
Paul Fultz II's avatar
Paul Fultz II committed
130
    compile_gen.cpp
Shucai Xiao's avatar
Shucai Xiao committed
131
    compile_hip.cpp
132
    compile_hip_code_object.cpp
133
    compile_miopen.cpp
134
    compiler.cpp
135
    device_name.cpp
136
    fuse_ck.cpp
Paul Fultz II's avatar
Paul Fultz II committed
137
    fuse_mlir.cpp
Paul's avatar
Paul committed
138
    fuse_ops.cpp
Shucai Xiao's avatar
Shucai Xiao committed
139
140
    gather.cpp
    gemm_impl.cpp
Paul's avatar
Paul committed
141
    hip.cpp
Shucai Xiao's avatar
Shucai Xiao committed
142
143
144
    int8_conv_pack.cpp
    int8_gemm_pack.cpp
    kernel.cpp
Paul's avatar
Paul committed
145
    lowering.cpp
146
    logsoftmax.cpp
Shucai Xiao's avatar
Shucai Xiao committed
147
    loop.cpp
Khalique's avatar
Khalique committed
148
    lrn.cpp
Paul Fultz II's avatar
Paul Fultz II committed
149
    mlir.cpp
turneram's avatar
turneram committed
150
    multinomial.cpp
151
    no_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
152
    nonzero.cpp
Shucai Xiao's avatar
Shucai Xiao committed
153
    pack_args.cpp
154
    pack_int8_args.cpp
155
    prefuse_ops.cpp
Shucai Xiao's avatar
Shucai Xiao committed
156
    pad.cpp
157
    perfdb.cpp
Shucai Xiao's avatar
Shucai Xiao committed
158
    pooling.cpp
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
159
    reverse.cpp
Shucai Xiao's avatar
Shucai Xiao committed
160
    rnn_variable_seq_lens.cpp
Shucai Xiao's avatar
Shucai Xiao committed
161
    rocblas.cpp
162
    scatter.cpp
Shucai Xiao's avatar
Shucai Xiao committed
163
    schedule_model.cpp
164
    sync_device.cpp
Shucai Xiao's avatar
Shucai Xiao committed
165
    target.cpp
166
    time_op.cpp
Shucai Xiao's avatar
Shucai Xiao committed
167
    topk.cpp
Shucai Xiao's avatar
Shucai Xiao committed
168
    write_literals.cpp
169
    ${JIT_GPU_SRCS}
Paul's avatar
Paul committed
170
)
171

Paul's avatar
Paul committed
172
set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
173
migraphx_generate_export_header(migraphx_gpu)
174

Artur Wojcik's avatar
Artur Wojcik committed
175
176
177
178
179
180
181
182
183
184
185
186
187
if(WIN32)
    target_include_directories(migraphx_gpu PRIVATE kernels)
    target_sources(migraphx_gpu
        PRIVATE
            kernels/migraphx_kernels.cpp
            kernels/migraphx_kernels.hpp
            kernels/resource.h
            kernels/resource.rc)
    target_compile_options(migraphx_gpu PUBLIC $<BUILD_INTERFACE:-Wno-ignored-attributes>)
else()
    target_link_libraries(migraphx_gpu PRIVATE migraphx_kernels)
    target_link_options(migraphx_gpu PUBLIC -Wno-option-ignored)
endif()
Artur Wojcik's avatar
Artur Wojcik committed
188

189
190
191
192
193
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
194

195
196
197
198
199
register_migraphx_gpu_ops(hip_
    argmax
    argmin
    gather
    logsoftmax
Shucai Xiao's avatar
Shucai Xiao committed
200
    loop
turneram's avatar
turneram committed
201
    multinomial
Shucai Xiao's avatar
Shucai Xiao committed
202
    nonzero
203
    pad
204
    prefix_scan_sum
Cagri Eryilmaz's avatar
Cagri Eryilmaz committed
205
    reverse
206
    scatter
Shucai Xiao's avatar
Shucai Xiao committed
207
    topk
208
209
210
211
212
213
214
215
)
register_migraphx_gpu_ops(miopen_
    abs
    contiguous
    int8_conv_pack
    lrn
    pooling
)
Umang Yadav's avatar
Umang Yadav committed
216
217
register_op(migraphx_gpu
    HEADER migraphx/gpu/rnn_variable_seq_lens.hpp
218
219
    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
220
221
register_op(migraphx_gpu
    HEADER migraphx/gpu/int8_gemm_pack.hpp
222
223
    OPERATORS gpu::hip_int8_gemm_pack_a gpu::hip_int8_gemm_pack_b
    INCLUDES migraphx/gpu/context.hpp)
Umang Yadav's avatar
Umang Yadav committed
224
225
register_op(migraphx_gpu
    HEADER migraphx/gpu/gemm.hpp
226
227
    OPERATORS gpu::rocblas_gemm<op::dot> gpu::rocblas_gemm<op::quant_dot>
    INCLUDES migraphx/gpu/context.hpp)
Umang Yadav's avatar
Umang Yadav committed
228
register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
229
    OPERATORS gpu::miopen_convolution<op::convolution> gpu::miopen_convolution<op::convolution_backwards> gpu::miopen_convolution<op::quant_convolution>
230
    INCLUDES migraphx/gpu/context.hpp)
Paul's avatar
Paul committed
231
rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
Paul's avatar
Paul committed
232
rocm_clang_tidy_check(migraphx_gpu)
Paul Fultz II's avatar
Paul Fultz II committed
233

234
set(MIGRAPHX_ENABLE_MLIR OFF CACHE BOOL "")
Umang Yadav's avatar
Umang Yadav committed
235

236
if(MIGRAPHX_ENABLE_MLIR)
237
238
239
    # 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
240
    target_compile_definitions(migraphx_gpu PRIVATE "-DMIGRAPHX_MLIR")
241
242
243
    # 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)
244
245
endif()

246
if(MIGRAPHX_USE_HIPRTC)
Umang Yadav's avatar
Umang Yadav committed
247
248
    message(STATUS "MIGraphX is using hipRTC")
    target_compile_definitions(migraphx_gpu PRIVATE -DMIGRAPHX_USE_HIPRTC=1)
249
else()
Umang Yadav's avatar
Umang Yadav committed
250
    message(STATUS "MIGraphX is using HIP Clang")
251

Umang Yadav's avatar
Umang Yadav committed
252
253
254
255
256
257
258
259
260
261
    # Get flags needed to compile hip
    include(TargetFlags)
    target_flags(HIP_COMPILER_FLAGS hip::device)

    # Remove cuda arch flags
    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}")

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

Umang Yadav's avatar
Umang Yadav committed
263
264
265
266
    foreach(_unused RANGE 2)
        string(REGEX REPLACE " /[^ ]+\\.(a|so) " " " HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
    endforeach()

Artur Wojcik's avatar
quots  
Artur Wojcik committed
267
    message(STATUS "Hip compiler flags: \"${HIP_COMPILER_FLAGS}\"")
Umang Yadav's avatar
Umang Yadav committed
268
    target_compile_definitions(migraphx_gpu PRIVATE
Artur Wojcik's avatar
quots  
Artur Wojcik committed
269
270
        -DMIGRAPHX_HIP_COMPILER="${CMAKE_CXX_COMPILER}"
        -DMIGRAPHX_HIP_COMPILER_FLAGS="${HIP_COMPILER_FLAGS}"
Umang Yadav's avatar
Umang Yadav committed
271
272
273
274
275
    )

    if(DEFINED CMAKE_CXX_COMPILER_LAUNCHER)
        execute_process(COMMAND which ${CMAKE_CXX_COMPILER_LAUNCHER} OUTPUT_VARIABLE MIGRAPHX_HIP_COMPILER_LAUNCHER)
        string(STRIP "${MIGRAPHX_HIP_COMPILER_LAUNCHER}" MIGRAPHX_HIP_COMPILER_LAUNCHER)
Artur Wojcik's avatar
quots  
Artur Wojcik committed
276
        target_compile_definitions(migraphx_gpu PRIVATE -DMIGRAPHX_HIP_COMPILER_LAUNCHER="${MIGRAPHX_HIP_COMPILER_LAUNCHER}")
Umang Yadav's avatar
Umang Yadav committed
277
    endif()
278
endif()
Paul Fultz II's avatar
Paul Fultz II committed
279

280
281
282
283
# Check miopen find mode api
include(CheckLibraryExists)
get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
284
check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
285

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

Umang Yadav's avatar
Umang Yadav committed
288
if(MIGRAPHX_USE_FIND_2_API)
289
290
291
292
293
294
    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()
295
296
297
298
    message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
else()
    message(STATUS "MIGraphx is using legacy Find API in MIOpen")
endif()
299

300
301
if(HAS_FIND_MODE_API)
    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
302
    message(STATUS "MIGraphx is using Find Mode API of MIOpen")
303
304
305
306
else()
    message(STATUS "MIOpen does not have find mode api")
endif()

Paul's avatar
Paul committed
307
target_link_libraries(migraphx_gpu PUBLIC migraphx MIOpen roc::rocblas)
Artur Wojcik's avatar
Artur Wojcik committed
308
target_link_libraries(migraphx_gpu PRIVATE migraphx_device)
309
310
311
312
if(NOT WIN32)
    # TODO: re-enable when CK is ported to Windows
    target_link_libraries(migraphx_gpu PRIVATE composable_kernel::jit_library)
endif()
mei-ye's avatar
mei-ye committed
313

314
add_subdirectory(driver)
315
add_subdirectory(hiprtc)
316

Paul's avatar
Paul committed
317
rocm_install_targets(
Umang Yadav's avatar
Umang Yadav committed
318
319
    TARGETS migraphx_gpu migraphx_device compile_for_gpu
    INCLUDE
Paul's avatar
Paul committed
320
321
    ${CMAKE_CURRENT_SOURCE_DIR}/include
)