CMakeLists.txt 15.2 KB
Newer Older
1
cmake_minimum_required(VERSION 3.14)
JD's avatar
JD committed
2
3

# Check support for CUDA/HIP in Cmake
Chao Liu's avatar
Chao Liu committed
4
project(composable_kernel)
Chao Liu's avatar
Chao Liu committed
5

6
7
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")

8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
if (DTYPES)
        add_definitions(-DDTYPES)
        if (DTYPES MATCHES "int8")
                add_definitions(-D__int8__)
        endif()
        if (DTYPES MATCHES "fp8")
                add_definitions(-D__fp8__)
        endif()
        if (DTYPES MATCHES "fp16")
                add_definitions(-D__fp16__)
        endif()
        if (DTYPES MATCHES "fp32")
                add_definitions(-D__fp32__)
        endif()
        if (DTYPES MATCHES "fp64")
                add_definitions(-D__fp64__)
        endif()
        if (DTYPES MATCHES "bf16")
                add_definitions(-D__bf16__)
        endif()
        message("DTYPES macro set to ${DTYPES}")
else()
        add_definitions(-D__int8__ -D__fp8__ -D__fp16__ -D__fp32__ -D__fp64__ -D__bf16__)
endif()

33
34
35
36
37
38
39
40
if(DL_KERNELS)
    add_definitions(-DDL_KERNELS)
endif()

if(INSTANCES_ONLY)
    add_definitions(-DINSTANCES_ONLY)
endif()

JD's avatar
JD committed
41
42
enable_testing()

43
set(ROCM_SYMLINK_LIBS OFF)
Anthony Chang's avatar
Anthony Chang committed
44
find_package(ROCM REQUIRED PATHS /opt/rocm)
JD's avatar
JD committed
45
46
47
48
49
50

include(ROCMInstallTargets)
include(ROCMPackageConfigHelpers)
include(ROCMSetupVersion)
include(ROCMInstallSymlinks)
include(ROCMCreatePackage)
Chao Liu's avatar
Chao Liu committed
51
include(CheckCXXCompilerFlag)
52
include(ROCMCheckTargetIds)
53
rocm_setup_version(VERSION 0.2.0)
JD's avatar
JD committed
54
55
56
include(TargetFlags)
list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX} ${CMAKE_INSTALL_PREFIX}/llvm ${CMAKE_INSTALL_PREFIX}/hip /opt/rocm /opt/rocm/llvm /opt/rocm/hip)

57
58
59
60
61
62
message("GPU_TARGETS= ${GPU_TARGETS}")

message("checking which targets are supported")
#This is the list of targets to be used in case GPU_TARGETS is not set on command line
#These targets will be filtered and only supported ones will be used
#Setting GPU_TARGETS on command line will override this list
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
if(NOT PROFILER_ONLY)
    rocm_check_target_ids(DEFAULT_GPU_TARGETS
        TARGETS "gfx900;gfx906;gfx908;gfx90a;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102")
else()
    add_definitions(-DPROFILER_ONLY)
    if(GPU_TARGETS)
        message(FATAL_ERROR "For PROFILE_ONLY build, please do not set GPU_TARGETS, use GPU_ARCH = gfx9, gfx10, or gfx11")
    endif()
    if(GPU_ARCH MATCHES "gfx9")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx900;gfx906;gfx908;gfx90a;gfx940;gfx941;gfx942")
    elseif(GPU_ARCH MATCHES "gfx10")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx1030")
    elseif(GPU_ARCH MATCHES "gfx11")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx1100;gfx1101;gfx1102")
    else()
        message(FATAL_ERROR "For PROFILE_ONLY build, please specify GPU_ARCH as gfx9, gfx10, or gfx11")
    endif()
endif()

82
message("Supported GPU_TARGETS= ${DEFAULT_GPU_TARGETS}")
83

84
set(AMDGPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " ")
85
86
87
88
89
90

if(GPU_TARGETS)
    message("Building CK for the following targets: ${GPU_TARGETS}")
else()
    message("Building CK for the following targets: ${AMDGPU_TARGETS}")
endif()
91
find_package(hip)
92
93
94
95
# No assumption that HIP kernels are launched with uniform block size for backward compatibility
# SWDEV-413293 and https://reviews.llvm.org/D155213
math(EXPR hip_VERSION_FLAT "(${hip_VERSION_MAJOR} * 1000 + ${hip_VERSION_MINOR}) * 100000 + ${hip_VERSION_PATCH}")
message("hip_version_flat=${hip_VERSION_FLAT}")
96
if(${hip_VERSION_FLAT} GREATER 500723302)
97
98
99
   message("Adding the fno-offload-uniform-block compiler flag")
   add_compile_options(-fno-offload-uniform-block)
endif()
100

Adam Osewski's avatar
Adam Osewski committed
101
option(USE_BITINT_EXTENSION_INT4, "Whether to enable clang's BitInt extension to provide int4 data type." OFF)
102
option(USE_OPT_NAVI3X, "Whether to enable LDS cumode and Wavefront32 mode for NAVI3X silicons." OFF)
Adam Osewski's avatar
Adam Osewski committed
103
104
105
106
107
108
109

if(USE_BITINT_EXTENSION_INT4)
    add_compile_definitions(CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4)
    add_compile_options(-Wno-bit-int-extension)
    message("CK compiled with USE_BITINT_EXTENSION_INT4 set to ${USE_BITINT_EXTENSION_INT4}")
endif()

110
111
112
113
114
115
if(USE_OPT_NAVI3X)
    add_compile_options(-mcumode)
    add_compile_options(-mno-wavefrontsize64)
    message("CK compiled with USE_OPT_NAVI3X set to ${USE_OPT_NAVI3X}")
endif()

116
117
118
119
120
## Threads
set(THREADS_PREFER_PTHREAD_FLAG ON)
find_package(Threads REQUIRED)
link_libraries(Threads::Threads)

121
## C++
Chao Liu's avatar
Chao Liu committed
122
enable_language(CXX)
Chao Liu's avatar
Chao Liu committed
123
set(CMAKE_CXX_STANDARD 17)
Chao Liu's avatar
Chao Liu committed
124
set(CMAKE_CXX_STANDARD_REQUIRED ON)
Chao Liu's avatar
Chao Liu committed
125
set(CMAKE_CXX_EXTENSIONS OFF)
Chao Liu's avatar
Chao Liu committed
126
127
message("CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")

128
## OpenMP
Chao Liu's avatar
Chao Liu committed
129
130
131
132
133
134
135
136
137
138
139
if(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
	# workaround issue hipcc in rocm3.5 cannot find openmp
	set(OpenMP_CXX "${CMAKE_CXX_COMPILER}")
	set(OpenMP_CXX_FLAGS "-fopenmp=libomp -Wno-unused-command-line-argument")
	set(OpenMP_CXX_LIB_NAMES "libomp" "libgomp" "libiomp5")
	set(OpenMP_libomp_LIBRARY ${OpenMP_CXX_LIB_NAMES})
	set(OpenMP_libgomp_LIBRARY ${OpenMP_CXX_LIB_NAMES})
	set(OpenMP_libiomp5_LIBRARY ${OpenMP_CXX_LIB_NAMES})
else()
	find_package(OpenMP REQUIRED)
endif()
Chao Liu's avatar
Chao Liu committed
140

Chao Liu's avatar
Chao Liu committed
141
142
143
144
message("OpenMP_CXX_LIB_NAMES: ${OpenMP_CXX_LIB_NAMES}")
message("OpenMP_gomp_LIBRARY: ${OpenMP_gomp_LIBRARY}")
message("OpenMP_pthread_LIBRARY: ${OpenMP_pthread_LIBRARY}")
message("OpenMP_CXX_FLAGS: ${OpenMP_CXX_FLAGS}")
Chao Liu's avatar
Chao Liu committed
145

Chao Liu's avatar
Chao Liu committed
146
147
link_libraries(${OpenMP_gomp_LIBRARY})
link_libraries(${OpenMP_pthread_LIBRARY})
Chao Liu's avatar
Chao Liu committed
148

149
150
## HIP
find_package(HIP REQUIRED)
JD's avatar
JD committed
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
# Override HIP version in config.h, if necessary.
# The variables set by find_package() can't be overwritten,
# therefore let's use intermediate variables.
set(CK_HIP_VERSION_MAJOR "${HIP_VERSION_MAJOR}")
set(CK_HIP_VERSION_MINOR "${HIP_VERSION_MINOR}")
set(CK_HIP_VERSION_PATCH "${HIP_VERSION_PATCH}")
if( DEFINED CK_OVERRIDE_HIP_VERSION_MAJOR )
    set(CK_HIP_VERSION_MAJOR "${CK_OVERRIDE_HIP_VERSION_MAJOR}")
    message(STATUS "CK_HIP_VERSION_MAJOR overriden with ${CK_OVERRIDE_HIP_VERSION_MAJOR}")
endif()
if( DEFINED CK_OVERRIDE_HIP_VERSION_MINOR )
    set(CK_HIP_VERSION_MINOR "${CK_OVERRIDE_HIP_VERSION_MINOR}")
    message(STATUS "CK_HIP_VERSION_MINOR overriden with ${CK_OVERRIDE_HIP_VERSION_MINOR}")
endif()
if( DEFINED CK_OVERRIDE_HIP_VERSION_PATCH )
    set(CK_HIP_VERSION_PATCH "${CK_OVERRIDE_HIP_VERSION_PATCH}")
    message(STATUS "CK_HIP_VERSION_PATCH overriden with ${CK_OVERRIDE_HIP_VERSION_PATCH}")
endif()
message(STATUS "Build with HIP ${HIP_VERSION}")
170
171
link_libraries(hip::device)
add_compile_definitions(__HIP_PLATFORM_HCC__=1)
172

Chao Liu's avatar
Chao Liu committed
173
174
## tidy
include(EnableCompilerWarnings)
JD's avatar
JD committed
175
set(CK_TIDY_ERRORS ERRORS * -readability-inconsistent-declaration-parameter-name)
Chao Liu's avatar
Chao Liu committed
176
if(CMAKE_CXX_COMPILER MATCHES ".*hcc" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
JD's avatar
JD committed
177
    set(CK_TIDY_CHECKS -modernize-use-override -readability-non-const-parameter)
Chao Liu's avatar
Chao Liu committed
178
# Enable tidy on hip
JD's avatar
JD committed
179
180
elseif(CK_BACKEND STREQUAL "HIP" OR CK_BACKEND STREQUAL "HIPNOGPU")
    set(CK_TIDY_ERRORS ALL)
Chao Liu's avatar
Chao Liu committed
181
182
endif()

JD's avatar
JD committed
183

Chao Liu's avatar
Chao Liu committed
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
include(ClangTidy)
enable_clang_tidy(
    CHECKS
        *
        -abseil-*
        -android-cloexec-fopen
        # Yea we shouldn't be using rand()
        -cert-msc30-c
        -bugprone-exception-escape
        -bugprone-macro-parentheses
        -cert-env33-c
        -cert-msc32-c
        -cert-msc50-cpp
        -cert-msc51-cpp
        -cert-dcl37-c
        -cert-dcl51-cpp
        -clang-analyzer-alpha.core.CastToStruct
        -clang-analyzer-optin.performance.Padding
        -clang-diagnostic-deprecated-declarations
        -clang-diagnostic-extern-c-compat
        -clang-diagnostic-unused-command-line-argument
        -cppcoreguidelines-avoid-c-arrays
        -cppcoreguidelines-avoid-magic-numbers
        -cppcoreguidelines-explicit-virtual-functions
        -cppcoreguidelines-init-variables
        -cppcoreguidelines-macro-usage
        -cppcoreguidelines-non-private-member-variables-in-classes
        -cppcoreguidelines-pro-bounds-array-to-pointer-decay
        -cppcoreguidelines-pro-bounds-constant-array-index
        -cppcoreguidelines-pro-bounds-pointer-arithmetic
        -cppcoreguidelines-pro-type-member-init
        -cppcoreguidelines-pro-type-reinterpret-cast
        -cppcoreguidelines-pro-type-union-access
        -cppcoreguidelines-pro-type-vararg
        -cppcoreguidelines-special-member-functions
        -fuchsia-*
        -google-explicit-constructor
        -google-readability-braces-around-statements
        -google-readability-todo
        -google-runtime-int
        -google-runtime-references
        -hicpp-vararg
        -hicpp-braces-around-statements
        -hicpp-explicit-conversions
        -hicpp-named-parameter
        -hicpp-no-array-decay
        # We really shouldn't use bitwise operators with signed integers, but
        # opencl leaves us no choice
        -hicpp-avoid-c-arrays
        -hicpp-signed-bitwise
        -hicpp-special-member-functions
        -hicpp-uppercase-literal-suffix
        -hicpp-use-auto
        -hicpp-use-equals-default
        -hicpp-use-override
        -llvm-header-guard
        -llvm-include-order
        #-llvmlibc-*
        -llvmlibc-restrict-system-libc-headers
        -llvmlibc-callee-namespace
        -llvmlibc-implementation-in-namespace
        -llvm-else-after-return
        -llvm-qualified-auto
        -misc-misplaced-const
        -misc-non-private-member-variables-in-classes
        -misc-no-recursion
        -modernize-avoid-bind
        -modernize-avoid-c-arrays
        -modernize-pass-by-value
        -modernize-use-auto
        -modernize-use-default-member-init
        -modernize-use-equals-default
        -modernize-use-trailing-return-type
        -modernize-use-transparent-functors
        -performance-unnecessary-value-param
        -readability-braces-around-statements
        -readability-else-after-return
        # we are not ready to use it, but very useful
        -readability-function-cognitive-complexity
        -readability-isolate-declaration
        -readability-magic-numbers
        -readability-named-parameter
        -readability-uppercase-literal-suffix
        -readability-convert-member-functions-to-static
        -readability-qualified-auto
        -readability-redundant-string-init
        # too many narrowing conversions in our code
        -bugprone-narrowing-conversions
        -cppcoreguidelines-narrowing-conversions
        -altera-struct-pack-align
        -cppcoreguidelines-prefer-member-initializer
JD's avatar
JD committed
275
276
        ${CK_TIDY_CHECKS}
        ${CK_TIDY_ERRORS}
Chao Liu's avatar
Chao Liu committed
277
278
279
    HEADER_FILTER
        "\.hpp$"
    EXTRA_ARGS
JD's avatar
JD committed
280
        -DCK_USE_CLANG_TIDY
Chao Liu's avatar
Chao Liu committed
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
)

include(CppCheck)
enable_cppcheck(
    CHECKS
        warning
        style
        performance
        portability
    SUPPRESS
        ConfigurationNotChecked
        constStatement
        duplicateCondition
        noExplicitConstructor
        passedByValue
Chao Liu's avatar
Chao Liu committed
296
        preprocessorErrorDirective
Chao Liu's avatar
Chao Liu committed
297
298
299
300
301
302
303
        shadowVariable
        unusedFunction
        unusedPrivateFunction
        unusedStructMember
        unmatchedSuppression
    FORCE
    SOURCES
Chao Liu's avatar
Chao Liu committed
304
        library/src
Chao Liu's avatar
Chao Liu committed
305
306
307
    INCLUDE
        ${CMAKE_CURRENT_SOURCE_DIR}/include
        ${CMAKE_CURRENT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
308
        ${CMAKE_CURRENT_SOURCE_DIR}/library/include
Chao Liu's avatar
Chao Liu committed
309
310
311
312
    DEFINE
        CPPCHECK=1
        __linux__=1
)
Chao Liu's avatar
Chao Liu committed
313

JD's avatar
JD committed
314
315
316
317
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin)

Chao Liu's avatar
Chao Liu committed
318
319
320
include_directories(BEFORE
    ${PROJECT_SOURCE_DIR}/include
    ${PROJECT_SOURCE_DIR}/library/include
321
    ${HIP_INCLUDE_DIRS}
JD's avatar
JD committed
322
)
Chao Liu's avatar
Chao Liu committed
323

324

JD's avatar
JD committed
325
326
SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV")
if(BUILD_DEV)
Chao Liu's avatar
Chao Liu committed
327
328
    add_compile_options(-Werror)
    add_compile_options(-Weverything)
JD's avatar
JD committed
329
330
331
endif()
message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}")

Anthony Chang's avatar
Anthony Chang committed
332
333
add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR})

334
335
336
337
file(GLOB_RECURSE INSTANCE_FILES "${PROJECT_SOURCE_DIR}/*/device_*_instance.cpp")
file(GLOB dir_list RELATIVE ${PROJECT_SOURCE_DIR}/library/src/tensor_operation_instance/gpu ${PROJECT_SOURCE_DIR}/library/src/tensor_operation_instance/gpu/*)
set(CK_DEVICE_INSTANCES)
FOREACH(subdir_path ${dir_list})
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
set(target_dir)
IF(IS_DIRECTORY "${PROJECT_SOURCE_DIR}/library/src/tensor_operation_instance/gpu/${subdir_path}")
    set(cmake_instance)
    file(READ "${PROJECT_SOURCE_DIR}/library/src/tensor_operation_instance/gpu/${subdir_path}/CMakeLists.txt" cmake_instance)
    set(add_inst 0)
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"fp8\" " AND DTYPES MATCHES "fp8")
            #message("fp8 instance found!")
            set(add_inst 1)
    endif()
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"fp16\"" AND DTYPES MATCHES "fp16")
            #message("fp16 instance found!")
            set(add_inst 1)
    endif()
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"fp32\"" AND DTYPES MATCHES "fp32")
            #message("fp32 instance found!")
            set(add_inst 1)
    endif()
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"fp64\"" AND DTYPES MATCHES "fp64")
            #message("fp64 instance found!")
            set(add_inst 1)
    endif()
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"bf16\"" AND DTYPES MATCHES "bf16")
            #message("bf16 instance found!")
            set(add_inst 1)
    endif()
    if("${cmake_instance}" MATCHES "DTYPES MATCHES \"int8\"" AND DTYPES MATCHES "int8")
            #message("int8 instance found!")
            set(add_inst 1)
    endif()
    if(NOT "${cmake_instance}" MATCHES "DTYPES")
            #message("instance should be built for all types!")
            set(add_inst 1)
    endif()
    if(add_inst EQUAL 1 OR NOT DEFINED DTYPES)
      list(APPEND CK_DEVICE_INSTANCES device_${subdir_path}_instance)
    endif()
ENDIF()
375
ENDFOREACH()
376

377
add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES}  SOURCES ${INSTANCE_FILES})
378
add_subdirectory(library)
379

380
if(NOT DEFINED INSTANCES_ONLY)
381
 if(NOT DEFINED PROFILER_ONLY)
382
   rocm_package_setup_component(tests
383
384
        LIBRARY_NAME composablekernel
        PACKAGE_NAME tests # Prevent -static suffix on package name
385
   )
386

387
   rocm_package_setup_component(examples
388
389
        LIBRARY_NAME composablekernel
        PACKAGE_NAME examples
390
   )
391
392
   add_subdirectory(example)
   add_subdirectory(test)
393

394
   rocm_package_setup_component(profiler
395
396
        LIBRARY_NAME composablekernel
        PACKAGE_NAME ckProfiler
397
398
   )
   add_subdirectory(profiler)
399
400
401
402
403
404
405
406
  else()
    #When building PROFILER_ONLY, label the package with GPU_ARCH
    rocm_package_setup_component(profiler
       LIBRARY_NAME composablekernel
       PACKAGE_NAME ckProfiler_${GPU_ARCH}
    )
    add_subdirectory(profiler)
  endif()
407
endif()
JD's avatar
JD committed
408
409
410
411
412
413
414
415
416
417
418

#Create an interface target for the include only files and call it "composablekernels"
include(CMakePackageConfigHelpers)

set(version 1.0.0)
write_basic_package_version_file(
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake"
    VERSION "${version}"
    COMPATIBILITY AnyNewerVersion
)

Anthony Chang's avatar
Anthony Chang committed
419
configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in
JD's avatar
JD committed
420
        "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
Anthony Chang's avatar
Anthony Chang committed
421
        INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
JD's avatar
JD committed
422
423
424
        NO_CHECK_REQUIRED_COMPONENTS_MACRO
)

425
rocm_install(FILES
JD's avatar
JD committed
426
427
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake"
Anthony Chang's avatar
Anthony Chang committed
428
    DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
JD's avatar
JD committed
429
)
430
431
432
433
434
435
436
437
438
439
440

set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE")
set(CPACK_RPM_PACKAGE_LICENSE "MIT")

rocm_create_package(
    NAME composablekernel
    DESCRIPTION "High Performance Composable Kernel for AMD GPUs"
    MAINTAINER "MIOpen Kernels Dev Team <dl.MIOpen@amd.com>"
    LDCONFIG
    HEADER_ONLY
)