CMakeLists.txt 22.1 KB
Newer Older
1
cmake_minimum_required(VERSION 3.14)
2
3
4
5
if(POLICY CMP0140)
  # policies CMP0140 not known to CMake until 3.25
  cmake_policy(SET CMP0140 NEW)
endif()
JD's avatar
JD committed
6

7
8
get_property(_GENERATOR_IS_MULTI_CONFIG GLOBAL PROPERTY GENERATOR_IS_MULTI_CONFIG)

9
10
# This has to be initialized before the project() command appears
# Set the default of CMAKE_BUILD_TYPE to be release, unless user specifies with -D.  MSVC_IDE does not use CMAKE_BUILD_TYPE
11
12
13
14
15
16
if(_GENERATOR_IS_MULTI_CONFIG)
    set(CMAKE_CONFIGURATION_TYPES "Debug;Release;RelWithDebInfo;MinSizeRel" CACHE STRING
            "Available build types (configurations) on multi-config generators")
else()
    set(CMAKE_BUILD_TYPE Release CACHE STRING
            "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel.")
17
18
19
endif()

# Default installation path
20
if(NOT WIN32)
21
22
23
    set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "")
endif()

24
set(version 1.1.0)
JD's avatar
JD committed
25
# Check support for CUDA/HIP in Cmake
26
project(composable_kernel VERSION ${version} LANGUAGES CXX HIP)
27
include(CTest)
Chao Liu's avatar
Chao Liu committed
28

29
find_package(Python3 3.6 COMPONENTS Interpreter REQUIRED)
carlushuang's avatar
carlushuang committed
30

31
32
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")

33
if (DTYPES)
34
35
36
37
38
39
40
41
    add_definitions(-DDTYPES)
    if (DTYPES MATCHES "int8")
        add_definitions(-DCK_ENABLE_INT8)
        set(CK_ENABLE_INT8 "ON")
    endif()
    if (DTYPES MATCHES "fp8")
        add_definitions(-DCK_ENABLE_FP8)
        set(CK_ENABLE_FP8 "ON")
42
43
44
45
    endif()
    if (DTYPES MATCHES "bf8")
        add_definitions(-DCK_ENABLE_BF8)
        set(CK_ENABLE_BF8 "ON")
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
    endif()
    if (DTYPES MATCHES "fp16")
        add_definitions(-DCK_ENABLE_FP16)
        set(CK_ENABLE_FP16 "ON")
    endif()
    if (DTYPES MATCHES "fp32")
        add_definitions(-DCK_ENABLE_FP32)
        set(CK_ENABLE_FP32 "ON")
    endif()
    if (DTYPES MATCHES "fp64")
        add_definitions(-DCK_ENABLE_FP64)
        set(CK_ENABLE_FP64 "ON")
    endif()
    if (DTYPES MATCHES "bf16")
        add_definitions(-DCK_ENABLE_BF16)
        set(CK_ENABLE_BF16 "ON")
    endif()
    message("DTYPES macro set to ${DTYPES}")
64
else()
65
66
67
68
69
70
71
72
73
74
75
    add_definitions(-DCK_ENABLE_INT8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_BF16)
    set(CK_ENABLE_INT8 "ON")
    set(CK_ENABLE_FP16 "ON")
    set(CK_ENABLE_FP32 "ON")
    set(CK_ENABLE_FP64 "ON")
    set(CK_ENABLE_BF16 "ON")
    if (GPU_TARGETS MATCHES "gfx94")
        add_definitions(-DCK_ENABLE_FP8 -DCK_ENABLE_BF8)
        set(CK_ENABLE_FP8 "ON")
        set(CK_ENABLE_BF8 "ON")
    endif()
76
77
endif()

78
79
#for f8/bf8_t type
add_compile_options(-Wno-bit-int-extension)
80
add_compile_options(-Wno-pass-failed)
81
add_compile_options(-Wno-switch-default)
82

83
84
if(DL_KERNELS)
    add_definitions(-DDL_KERNELS)
85
    set(CK_ENABLE_DL_KERNELS "ON")
86
87
88
89
endif()

if(INSTANCES_ONLY)
    add_definitions(-DINSTANCES_ONLY)
90
    set(CK_ENABLE_INSTANCES_ONLY "ON")
91
92
endif()

93
94
include(getopt)

95
96
97
# CK version file to record release version as well as git commit hash
find_package(Git REQUIRED)
execute_process(COMMAND "${GIT_EXECUTABLE}" rev-parse HEAD OUTPUT_VARIABLE COMMIT_ID OUTPUT_STRIP_TRAILING_WHITESPACE)
98
configure_file(include/ck/version.h.in ${CMAKE_CURRENT_BINARY_DIR}/include/ck/version.h)
JD's avatar
JD committed
99

100
set(ROCM_SYMLINK_LIBS OFF)
Anthony Chang's avatar
Anthony Chang committed
101
find_package(ROCM REQUIRED PATHS /opt/rocm)
JD's avatar
JD committed
102
103
104
105
106
107

include(ROCMInstallTargets)
include(ROCMPackageConfigHelpers)
include(ROCMSetupVersion)
include(ROCMInstallSymlinks)
include(ROCMCreatePackage)
Chao Liu's avatar
Chao Liu committed
108
include(CheckCXXCompilerFlag)
109
include(ROCMCheckTargetIds)
JD's avatar
JD committed
110
include(TargetFlags)
111
112
113

rocm_setup_version(VERSION ${version})

114
list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX} ${CMAKE_INSTALL_PREFIX}/llvm ${CMAKE_INSTALL_PREFIX}/hip /opt/rocm /opt/rocm/llvm /opt/rocm/hip "$ENV{ROCM_PATH}" "$ENV{HIP_PATH}")
JD's avatar
JD committed
115

116
117
message("GPU_TARGETS= ${GPU_TARGETS}")

118
119
120
121
122
123
find_package(hip)
# 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}")

124
125
126
127
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
128
if(NOT PROFILER_ONLY)
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
    if(NOT ENABLE_ASAN_PACKAGING)
        #build CK for all supported targets
        if(NOT WIN32 AND ${hip_VERSION_FLAT} LESS 600300000)
            # WORKAROUND: compiler does not yet fully support gfx12 targets, need to fix version above
            rocm_check_target_ids(DEFAULT_GPU_TARGETS
                TARGETS "gfx908;gfx90a;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102")
        else()
            rocm_check_target_ids(DEFAULT_GPU_TARGETS
                TARGETS "gfx908;gfx90a;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1200;gfx1201")
        endif()
    else()
        #build CK only for xnack-supported targets
        rocm_check_target_ids(DEFAULT_GPU_TARGETS
            TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx940:xnack+;gfx941:xnack+;gfx942:xnack+")
        set(GPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " " FORCE)
    endif()
145
146
else()
    add_definitions(-DPROFILER_ONLY)
147
    set(GPU_TARGETS "" CACHE STRING "" FORCE)
148
    if(GPU_TARGETS)
149
        message(FATAL_ERROR "For PROFILE_ONLY build, please do not set GPU_TARGETS, use GPU_ARCH = gfx90, gfx94, gfx10, gfx11 or gfx12")
150
    endif()
151
    if(GPU_ARCH MATCHES "gfx90")
152
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx908;gfx90a")
153
154
    elseif(GPU_ARCH MATCHES "gfx94")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx940;gfx941;gfx942")
155
156
157
158
    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")
159
160
    elseif(GPU_ARCH MATCHES "gfx12")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx1200;gfx1201")
161
    else()
162
        message(FATAL_ERROR "For PROFILE_ONLY build, please specify GPU_ARCH as gfx90, gfx94, gfx10, gfx11 or gfx12")
163
    endif()
164
    set(GPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " " FORCE)
165
166
endif()

167
message("Supported GPU_TARGETS= ${DEFAULT_GPU_TARGETS}")
168
169
170
171

if(GPU_TARGETS)
    message("Building CK for the following targets: ${GPU_TARGETS}")
else()
172
    message("Building CK for the default targets: ${DEFAULT_GPU_TARGETS}")
173
endif()
174
175
176
177
178
179

if (GPU_TARGETS)
    if (GPU_TARGETS MATCHES "gfx9")
        add_definitions(-DCK_USE_XDL)
        set(CK_USE_XDL "ON")
    endif()
180
    if (GPU_TARGETS MATCHES "gfx11" OR GPU_TARGETS MATCHES "gfx12")
181
182
183
184
185
186
187
188
189
        add_definitions(-DCK_USE_WMMA)
        set(CK_USE_WMMA "ON")
    endif()
else()
    add_definitions(-DCK_USE_WMMA -DCK_USE_XDL)
    set(CK_USE_XDL "ON")
    set(CK_USE_WMMA "ON")
endif()

190
191
192
# CK config file to record supported datatypes, etc.
configure_file(include/ck/config.h.in ${CMAKE_CURRENT_BINARY_DIR}/include/ck/config.h)

193
if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 500723302)
194
195
196
197
198
  check_cxx_compiler_flag("-fno-offload-uniform-block" HAS_NO_OFFLOAD_UNIFORM_BLOCK)
  if(HAS_NO_OFFLOAD_UNIFORM_BLOCK)
    message("Adding the fno-offload-uniform-block compiler flag")
    add_compile_options(-fno-offload-uniform-block)
  endif()
199
endif()
200
if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 600140090)
201
202
203
204
205
  check_cxx_compiler_flag("-mllvm -enable-post-misched=0" HAS_ENABLE_POST_MISCHED)
  if(HAS_ENABLE_POST_MISCHED)
    message("Adding the enable-post-misched=0 compiler flag")
    add_compile_options("SHELL: -mllvm -enable-post-misched=0")
  endif()
206
endif()
207
208
209
set(check-coerce)
check_cxx_compiler_flag(" -mllvm -amdgpu-coerce-illegal-types=1" check-coerce)
if(NOT WIN32 AND check-coerce AND ${hip_VERSION_FLAT} GREATER 600241132 AND ${hip_VERSION_FLAT} LESS 600300000)
210
211
212
213
214
215
216
   message("Adding the amdgpu-coerce-illegal-types=1")
   add_compile_options("SHELL: -mllvm -amdgpu-coerce-illegal-types=1")
endif()
if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 600241132)
   message("Adding -amdgpu-early-inline-all=true and -amdgpu-function-calls=false")
   add_compile_options("SHELL: -mllvm -amdgpu-early-inline-all=true")
   add_compile_options("SHELL: -mllvm -amdgpu-function-calls=false")
217
endif()
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
#
# Seperate linking jobs from compiling
# Too many concurrent linking jobs can break the build
# Copied from LLVM
set(CK_PARALLEL_LINK_JOBS "" CACHE STRING
  "Define the maximum number of concurrent link jobs (Ninja only).")
if(CMAKE_GENERATOR MATCHES "Ninja")
  if(CK_PARALLEL_LINK_JOBS)
    set_property(GLOBAL APPEND PROPERTY JOB_POOLS link_job_pool=${CK_PARALLEL_LINK_JOBS})
    set(CMAKE_JOB_POOL_LINK link_job_pool)
  endif()
elseif(CK_PARALLEL_LINK_JOBS)
  message(WARNING "Job pooling is only available with Ninja generators.")
endif()
# Similar for compiling
set(CK_PARALLEL_COMPILE_JOBS "" CACHE STRING
  "Define the maximum number of concurrent compile jobs (Ninja only).")
if(CMAKE_GENERATOR MATCHES "Ninja")
  if(CK_PARALLEL_COMPILE_JOBS)
    set_property(GLOBAL APPEND PROPERTY JOB_POOLS compile_job_pool=${CK_PARALLEL_COMPILE_JOBS})
    set(CMAKE_JOB_POOL_COMPILE compile_job_pool)
  endif()
elseif(CK_PARALLEL_COMPILE_JOBS)
  message(WARNING "Job pooling is only available with Ninja generators.")
endif()


245
option(USE_BITINT_EXTENSION_INT4 "Whether to enable clang's BitInt extension to provide int4 data type." OFF)
Illia Silin's avatar
Illia Silin committed
246
option(USE_OPT_GFX11 "Whether to enable LDS cumode and Wavefront32 mode for GFX11 silicons." OFF)
Adam Osewski's avatar
Adam Osewski committed
247
248
249
250
251
252
253

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()

Illia Silin's avatar
Illia Silin committed
254
if(USE_OPT_GFX11)
255
256
    add_compile_options(-mcumode)
    add_compile_options(-mno-wavefrontsize64)
Illia Silin's avatar
Illia Silin committed
257
    message("CK compiled with USE_OPT_GFX11 set to ${USE_OPT_GFX11}")
258
259
endif()

260
261
262
263
264
## Threads
set(THREADS_PREFER_PTHREAD_FLAG ON)
find_package(Threads REQUIRED)
link_libraries(Threads::Threads)

265
## C++
Chao Liu's avatar
Chao Liu committed
266
set(CMAKE_CXX_STANDARD 17)
Chao Liu's avatar
Chao Liu committed
267
set(CMAKE_CXX_STANDARD_REQUIRED ON)
Chao Liu's avatar
Chao Liu committed
268
set(CMAKE_CXX_EXTENSIONS OFF)
269
270
message("CMAKE_CXX_COMPILER: ${CMAKE_CXX_COMPILER}")

271
272
273
274
275
276
277
278
279
280
# https://gcc.gnu.org/onlinedocs/libstdc++/manual/using_macros.html
# _GLIBCXX_ASSERTIONS
# Undefined by default. When defined, enables extra error checking in the form of
# precondition assertions, such as bounds checking in strings and null pointer
# checks when dereferencing smart pointers
option(USE_GLIBCXX_ASSERTIONS "Turn on additional c++ library checks." OFF)
if(USE_GLIBCXX_ASSERTIONS)
  add_compile_options(-Wp,-D_GLIBCXX_ASSERTIONS)
endif()

281
282
283
284
285
## HIP
set(CMAKE_HIP_PLATFORM amd)
set(CMAKE_HIP_COMPILER ${CMAKE_CXX_COMPILER})
set(CMAKE_HIP_EXTENSIONS ON)
message("CMAKE_HIP_COMPILER: ${CMAKE_HIP_COMPILER}")
Chao Liu's avatar
Chao Liu committed
286

287
## OpenMP
Chao Liu's avatar
Chao Liu committed
288
289
290
291
292
293
294
295
296
297
298
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
299

Chao Liu's avatar
Chao Liu committed
300
301
302
303
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
304

Chao Liu's avatar
Chao Liu committed
305
306
link_libraries(${OpenMP_gomp_LIBRARY})
link_libraries(${OpenMP_pthread_LIBRARY})
Chao Liu's avatar
Chao Liu committed
307

308
309
## HIP
find_package(HIP REQUIRED)
JD's avatar
JD committed
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
# 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}")
329
link_libraries(hip::device)
330
331
332
333
334
if(CK_hip_VERSION VERSION_GREATER_EQUAL 6.0.23494)
    add_compile_definitions(__HIP_PLATFORM_AMD__=1)
else()
    add_compile_definitions(__HIP_PLATFORM_HCC__=1)
endif()
335

Chao Liu's avatar
Chao Liu committed
336
337
## tidy
include(EnableCompilerWarnings)
JD's avatar
JD committed
338
set(CK_TIDY_ERRORS ERRORS * -readability-inconsistent-declaration-parameter-name)
Chao Liu's avatar
Chao Liu committed
339
if(CMAKE_CXX_COMPILER MATCHES ".*hcc" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
JD's avatar
JD committed
340
    set(CK_TIDY_CHECKS -modernize-use-override -readability-non-const-parameter)
Chao Liu's avatar
Chao Liu committed
341
# Enable tidy on hip
JD's avatar
JD committed
342
343
elseif(CK_BACKEND STREQUAL "HIP" OR CK_BACKEND STREQUAL "HIPNOGPU")
    set(CK_TIDY_ERRORS ALL)
Chao Liu's avatar
Chao Liu committed
344
345
endif()

JD's avatar
JD committed
346

Chao Liu's avatar
Chao Liu committed
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
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
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
438
439
        ${CK_TIDY_CHECKS}
        ${CK_TIDY_ERRORS}
Chao Liu's avatar
Chao Liu committed
440
441
442
    HEADER_FILTER
        "\.hpp$"
    EXTRA_ARGS
JD's avatar
JD committed
443
        -DCK_USE_CLANG_TIDY
Chao Liu's avatar
Chao Liu committed
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
)

include(CppCheck)
enable_cppcheck(
    CHECKS
        warning
        style
        performance
        portability
    SUPPRESS
        ConfigurationNotChecked
        constStatement
        duplicateCondition
        noExplicitConstructor
        passedByValue
Chao Liu's avatar
Chao Liu committed
459
        preprocessorErrorDirective
Chao Liu's avatar
Chao Liu committed
460
461
462
463
464
465
466
        shadowVariable
        unusedFunction
        unusedPrivateFunction
        unusedStructMember
        unmatchedSuppression
    FORCE
    SOURCES
Chao Liu's avatar
Chao Liu committed
467
        library/src
Chao Liu's avatar
Chao Liu committed
468
469
470
    INCLUDE
        ${CMAKE_CURRENT_SOURCE_DIR}/include
        ${CMAKE_CURRENT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
471
        ${CMAKE_CURRENT_SOURCE_DIR}/library/include
Chao Liu's avatar
Chao Liu committed
472
473
474
475
    DEFINE
        CPPCHECK=1
        __linux__=1
)
Chao Liu's avatar
Chao Liu committed
476

JD's avatar
JD committed
477
478
479
480
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)

481
# set CK project include directories
Chao Liu's avatar
Chao Liu committed
482
include_directories(BEFORE
483
    ${PROJECT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
484
485
    ${PROJECT_SOURCE_DIR}/include
    ${PROJECT_SOURCE_DIR}/library/include
486
    ${HIP_INCLUDE_DIRS}
JD's avatar
JD committed
487
)
Chao Liu's avatar
Chao Liu committed
488

JD's avatar
JD committed
489
490
SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV")
if(BUILD_DEV)
491
492
    add_compile_options(-Werror)
    add_compile_options(-Weverything)
JD's avatar
JD committed
493
494
495
endif()
message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}")

496
497
498
499
500
501
502
if("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
    add_compile_options(-fcolor-diagnostics)
endif()
if("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU" AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 4.9)
    add_compile_options(-fdiagnostics-color=always)
endif()

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

505
506
507
508
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})
509
510
511
512
513
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)
514
    if(("${cmake_instance}" MATCHES "fp8" OR "${cmake_instance}" MATCHES "_f8") AND DTYPES MATCHES "fp8")
515
        set(add_inst 1)
516
    endif()
517
    if(("${cmake_instance}" MATCHES "bf8" OR "${cmake_instance}" MATCHES "_b8") AND DTYPES MATCHES "bf8")
518
519
        set(add_inst 1)
    endif()
520
    if(("${cmake_instance}" MATCHES "fp16" OR "${cmake_instance}" MATCHES "_f16") AND DTYPES MATCHES "fp16")
521
        set(add_inst 1)
522
    endif()
523
    if(("${cmake_instance}" MATCHES "fp32" OR "${cmake_instance}" MATCHES "_f32") AND DTYPES MATCHES "fp32")
524
        set(add_inst 1)
525
    endif()
526
    if(("${cmake_instance}" MATCHES "fp64" OR "${cmake_instance}" MATCHES "_f64") AND DTYPES MATCHES "fp64")
527
        set(add_inst 1)
528
    endif()
529
    if(("${cmake_instance}" MATCHES "bf16" OR "${cmake_instance}" MATCHES "_b16") AND DTYPES MATCHES "bf16")
530
        set(add_inst 1)
531
    endif()
532
    if(("${cmake_instance}" MATCHES "int8" OR "${cmake_instance}" MATCHES "_i8") AND DTYPES MATCHES "int8")
533
        set(add_inst 1)
534
535
    endif()
    if(NOT "${cmake_instance}" MATCHES "DTYPES")
536
        set(add_inst 1)
537
538
    endif()
    if(add_inst EQUAL 1 OR NOT DEFINED DTYPES)
539
        list(APPEND CK_DEVICE_INSTANCES device_${subdir_path}_instance)
540
541
    endif()
ENDIF()
542
ENDFOREACH()
543

544
add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES}  SOURCES ${INSTANCE_FILES})
545
add_subdirectory(library)
546

547
if(NOT DEFINED INSTANCES_ONLY)
548
 if(NOT DEFINED PROFILER_ONLY)
549
   rocm_package_setup_component(tests
550
551
        LIBRARY_NAME composablekernel
        PACKAGE_NAME tests # Prevent -static suffix on package name
552
   )
553

554
   rocm_package_setup_component(examples
555
556
        LIBRARY_NAME composablekernel
        PACKAGE_NAME examples
557
   )
558
   add_subdirectory(example)
arai713's avatar
arai713 committed
559
560
561
   if(GPU_TARGETS MATCHES "gfx9" AND NOT INSTANCES_ONLY)
      add_subdirectory(codegen)
   endif()
562
563
564
   if(BUILD_TESTING)
      add_subdirectory(test)
   endif()
565

566
   rocm_package_setup_component(profiler
567
        LIBRARY_NAME composablekernel
568
        PACKAGE_NAME ckprofiler
569
570
   )
   add_subdirectory(profiler)
571
572
573
574
  else()
    #When building PROFILER_ONLY, label the package with GPU_ARCH
    rocm_package_setup_component(profiler
       LIBRARY_NAME composablekernel
575
       PACKAGE_NAME ckprofiler_${GPU_ARCH}
576
577
578
    )
    add_subdirectory(profiler)
  endif()
579
endif()
JD's avatar
JD committed
580
581
582
583
584
585
586
587
588
589

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

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

Anthony Chang's avatar
Anthony Chang committed
590
configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in
591
592
593
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
    INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
    NO_CHECK_REQUIRED_COMPONENTS_MACRO
JD's avatar
JD committed
594
595
)

596
rocm_install(FILES
JD's avatar
JD committed
597
598
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake"
Anthony Chang's avatar
Anthony Chang committed
599
    DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
JD's avatar
JD committed
600
)
601

602
# Install CK version and configuration files
603
rocm_install(FILES
604
605
606
607
608
    ${PROJECT_BINARY_DIR}/include/ck/version.h
    ${PROJECT_BINARY_DIR}/include/ck/config.h
    DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/ck/
)

609
610
611
612
613
614
615
616
617
618
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
)