CMakeLists.txt 17.4 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

Artur Wojcik's avatar
Artur Wojcik committed
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
Artur Wojcik's avatar
Artur Wojcik committed
11
if(_GENERATOR_IS_MULTI_CONFIG)
Artur Wojcik's avatar
Artur Wojcik committed
12
13
    set(CMAKE_CONFIGURATION_TYPES "Debug;Release;RelWithDebInfo;MinSizeRel" CACHE STRING
            "Available build types (configurations) on multi-config generators")
Artur Wojcik's avatar
Artur Wojcik committed
14
else()
Artur Wojcik's avatar
Artur Wojcik committed
15
16
    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
Artur Wojcik's avatar
Artur Wojcik committed
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
Artur Wojcik's avatar
Artur Wojcik committed
26
27
project(composable_kernel VERSION ${version} LANGUAGES CXX)
include(CTest)
Chao Liu's avatar
Chao Liu committed
28

29
30
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")

31
if (DTYPES)
32
33
34
35
36
37
38
39
    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")
40
41
42
43
    endif()
    if (DTYPES MATCHES "bf8")
        add_definitions(-DCK_ENABLE_BF8)
        set(CK_ENABLE_BF8 "ON")
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
    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}")
62
else()
63
    add_definitions(-DCK_ENABLE_INT8 -DCK_ENABLE_FP8 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_BF16)
64
    set(CK_ENABLE_ALL_DTYPES "ON")
65
66
endif()

67
68
69
#for f8/bf8_t type
add_compile_options(-Wno-bit-int-extension)

70
71
if(DL_KERNELS)
    add_definitions(-DDL_KERNELS)
72
    set(CK_ENABLE_DL_KERNELS "ON")
73
74
75
76
endif()

if(INSTANCES_ONLY)
    add_definitions(-DINSTANCES_ONLY)
77
    set(CK_ENABLE_INSTANCES_ONLY "ON")
78
79
endif()

Artur Wojcik's avatar
Artur Wojcik committed
80
include(getopt)
Artur Wojcik's avatar
Artur Wojcik committed
81

82
# CK config file to record supported datatypes, etc.
Artur Wojcik's avatar
Artur Wojcik committed
83
configure_file(include/ck/config.h.in ${CMAKE_CURRENT_BINARY_DIR}/include/ck/config.h)
84
85
86
87

# 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)
Artur Wojcik's avatar
Artur Wojcik committed
88
configure_file(include/ck/version.h.in ${CMAKE_CURRENT_BINARY_DIR}/include/ck/version.h)
JD's avatar
JD committed
89

90
set(ROCM_SYMLINK_LIBS OFF)
Anthony Chang's avatar
Anthony Chang committed
91
find_package(ROCM REQUIRED PATHS /opt/rocm)
JD's avatar
JD committed
92
93
94
95
96
97

include(ROCMInstallTargets)
include(ROCMPackageConfigHelpers)
include(ROCMSetupVersion)
include(ROCMInstallSymlinks)
include(ROCMCreatePackage)
Chao Liu's avatar
Chao Liu committed
98
include(CheckCXXCompilerFlag)
99
include(ROCMCheckTargetIds)
JD's avatar
JD committed
100
include(TargetFlags)
101
102
103

rocm_setup_version(VERSION ${version})

Artur Wojcik's avatar
Artur Wojcik committed
104
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
105

106
107
108
109
110
111
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
112
113
if(NOT PROFILER_ONLY)
    rocm_check_target_ids(DEFAULT_GPU_TARGETS
114
        TARGETS "gfx908;gfx90a;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102")
115
116
else()
    add_definitions(-DPROFILER_ONLY)
117
    set(GPU_TARGETS "" CACHE STRING "" FORCE)
118
    if(GPU_TARGETS)
119
        message(FATAL_ERROR "For PROFILE_ONLY build, please do not set GPU_TARGETS, use GPU_ARCH = gfx90, gfx94, gfx10, or gfx11")
120
    endif()
121
    if(GPU_ARCH MATCHES "gfx90")
122
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx908;gfx90a")
123
124
    elseif(GPU_ARCH MATCHES "gfx94")
        rocm_check_target_ids(DEFAULT_GPU_TARGETS TARGETS "gfx940;gfx941;gfx942")
125
126
127
128
129
    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()
130
        message(FATAL_ERROR "For PROFILE_ONLY build, please specify GPU_ARCH as gfx90, gfx94, gfx10, or gfx11")
131
    endif()
132
    set(GPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " " FORCE)
133
134
endif()

135
message("Supported GPU_TARGETS= ${DEFAULT_GPU_TARGETS}")
136

137
set(AMDGPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " " FORCE)
138
139
140
141
142
143

if(GPU_TARGETS)
    message("Building CK for the following targets: ${GPU_TARGETS}")
else()
    message("Building CK for the following targets: ${AMDGPU_TARGETS}")
endif()
144
find_package(hip)
145
146
147
148
# 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}")
Artur Wojcik's avatar
Artur Wojcik committed
149
if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 500723302)
150
151
152
   message("Adding the fno-offload-uniform-block compiler flag")
   add_compile_options(-fno-offload-uniform-block)
endif()
153

Adam Osewski's avatar
Adam Osewski committed
154
option(USE_BITINT_EXTENSION_INT4, "Whether to enable clang's BitInt extension to provide int4 data type." OFF)
155
option(USE_OPT_NAVI3X, "Whether to enable LDS cumode and Wavefront32 mode for NAVI3X silicons." OFF)
Adam Osewski's avatar
Adam Osewski committed
156
157
158
159
160
161
162

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

163
164
165
166
167
168
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()

169
## Threads
Artur Wojcik's avatar
Artur Wojcik committed
170
set(THREADS_PREFER_PTHREAD_FLAG ON)
171
172
173
find_package(Threads REQUIRED)
link_libraries(Threads::Threads)

174
## C++
Chao Liu's avatar
Chao Liu committed
175
set(CMAKE_CXX_STANDARD 17)
Chao Liu's avatar
Chao Liu committed
176
set(CMAKE_CXX_STANDARD_REQUIRED ON)
Chao Liu's avatar
Chao Liu committed
177
set(CMAKE_CXX_EXTENSIONS OFF)
Chao Liu's avatar
Chao Liu committed
178
179
message("CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")

180
## OpenMP
Chao Liu's avatar
Chao Liu committed
181
182
183
184
185
186
187
188
189
190
191
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
192

Chao Liu's avatar
Chao Liu committed
193
194
195
196
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
197

Chao Liu's avatar
Chao Liu committed
198
199
link_libraries(${OpenMP_gomp_LIBRARY})
link_libraries(${OpenMP_pthread_LIBRARY})
Chao Liu's avatar
Chao Liu committed
200

201
202
## HIP
find_package(HIP REQUIRED)
JD's avatar
JD committed
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
# 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}")
222
223
link_libraries(hip::device)
add_compile_definitions(__HIP_PLATFORM_HCC__=1)
224

Chao Liu's avatar
Chao Liu committed
225
226
## tidy
include(EnableCompilerWarnings)
JD's avatar
JD committed
227
set(CK_TIDY_ERRORS ERRORS * -readability-inconsistent-declaration-parameter-name)
Chao Liu's avatar
Chao Liu committed
228
if(CMAKE_CXX_COMPILER MATCHES ".*hcc" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
JD's avatar
JD committed
229
    set(CK_TIDY_CHECKS -modernize-use-override -readability-non-const-parameter)
Chao Liu's avatar
Chao Liu committed
230
# Enable tidy on hip
JD's avatar
JD committed
231
232
elseif(CK_BACKEND STREQUAL "HIP" OR CK_BACKEND STREQUAL "HIPNOGPU")
    set(CK_TIDY_ERRORS ALL)
Chao Liu's avatar
Chao Liu committed
233
234
endif()

JD's avatar
JD committed
235

Chao Liu's avatar
Chao Liu committed
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
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
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
327
328
        ${CK_TIDY_CHECKS}
        ${CK_TIDY_ERRORS}
Chao Liu's avatar
Chao Liu committed
329
330
331
    HEADER_FILTER
        "\.hpp$"
    EXTRA_ARGS
JD's avatar
JD committed
332
        -DCK_USE_CLANG_TIDY
Chao Liu's avatar
Chao Liu committed
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
)

include(CppCheck)
enable_cppcheck(
    CHECKS
        warning
        style
        performance
        portability
    SUPPRESS
        ConfigurationNotChecked
        constStatement
        duplicateCondition
        noExplicitConstructor
        passedByValue
Chao Liu's avatar
Chao Liu committed
348
        preprocessorErrorDirective
Chao Liu's avatar
Chao Liu committed
349
350
351
352
353
354
355
        shadowVariable
        unusedFunction
        unusedPrivateFunction
        unusedStructMember
        unmatchedSuppression
    FORCE
    SOURCES
Chao Liu's avatar
Chao Liu committed
356
        library/src
Chao Liu's avatar
Chao Liu committed
357
358
359
    INCLUDE
        ${CMAKE_CURRENT_SOURCE_DIR}/include
        ${CMAKE_CURRENT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
360
        ${CMAKE_CURRENT_SOURCE_DIR}/library/include
Chao Liu's avatar
Chao Liu committed
361
362
363
364
    DEFINE
        CPPCHECK=1
        __linux__=1
)
Chao Liu's avatar
Chao Liu committed
365

JD's avatar
JD committed
366
367
368
369
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)

370
# set CK project include directories
Chao Liu's avatar
Chao Liu committed
371
include_directories(BEFORE
372
    ${PROJECT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
373
374
    ${PROJECT_SOURCE_DIR}/include
    ${PROJECT_SOURCE_DIR}/library/include
375
    ${HIP_INCLUDE_DIRS}
JD's avatar
JD committed
376
)
Chao Liu's avatar
Chao Liu committed
377

JD's avatar
JD committed
378
379
SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV")
if(BUILD_DEV)
380
381
    add_compile_options(-Werror)
    add_compile_options(-Weverything)
JD's avatar
JD committed
382
383
384
endif()
message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}")

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

387
388
389
390
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})
391
392
393
394
395
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)
396
    if(("${cmake_instance}" MATCHES "fp8" OR "${cmake_instance}" MATCHES "_f8") AND DTYPES MATCHES "fp8")
397
        set(add_inst 1)
398
    endif()
399
    if(("${cmake_instance}" MATCHES "bf8" OR "${cmake_instance}" MATCHES "_b8") AND DTYPES MATCHES "bf8")
400
401
        set(add_inst 1)
    endif()
402
    if(("${cmake_instance}" MATCHES "fp16" OR "${cmake_instance}" MATCHES "_f16") AND DTYPES MATCHES "fp16")
403
        set(add_inst 1)
404
    endif()
405
    if(("${cmake_instance}" MATCHES "fp32" OR "${cmake_instance}" MATCHES "_f32") AND DTYPES MATCHES "fp32")
406
        set(add_inst 1)
407
    endif()
408
    if(("${cmake_instance}" MATCHES "fp64" OR "${cmake_instance}" MATCHES "_f64") AND DTYPES MATCHES "fp64")
409
        set(add_inst 1)
410
    endif()
411
    if(("${cmake_instance}" MATCHES "bf16" OR "${cmake_instance}" MATCHES "_b16") AND DTYPES MATCHES "bf16")
412
        set(add_inst 1)
413
    endif()
414
    if(("${cmake_instance}" MATCHES "int8" OR "${cmake_instance}" MATCHES "_i8") AND DTYPES MATCHES "int8")
415
        set(add_inst 1)
416
417
    endif()
    if(NOT "${cmake_instance}" MATCHES "DTYPES")
418
        set(add_inst 1)
419
420
    endif()
    if(add_inst EQUAL 1 OR NOT DEFINED DTYPES)
421
        list(APPEND CK_DEVICE_INSTANCES device_${subdir_path}_instance)
422
423
    endif()
ENDIF()
424
ENDFOREACH()
425

426
add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES}  SOURCES ${INSTANCE_FILES})
427
add_subdirectory(library)
428

429
if(NOT DEFINED INSTANCES_ONLY)
430
 if(NOT DEFINED PROFILER_ONLY)
431
   rocm_package_setup_component(tests
432
433
        LIBRARY_NAME composablekernel
        PACKAGE_NAME tests # Prevent -static suffix on package name
434
   )
435

436
   rocm_package_setup_component(examples
437
438
        LIBRARY_NAME composablekernel
        PACKAGE_NAME examples
439
   )
440
   add_subdirectory(example)
Artur Wojcik's avatar
Artur Wojcik committed
441
442
443
   if(BUILD_TESTING)
     add_subdirectory(test)
   endif()
444

445
   rocm_package_setup_component(profiler
446
        LIBRARY_NAME composablekernel
447
        PACKAGE_NAME ckprofiler
448
449
   )
   add_subdirectory(profiler)
450
451
452
453
  else()
    #When building PROFILER_ONLY, label the package with GPU_ARCH
    rocm_package_setup_component(profiler
       LIBRARY_NAME composablekernel
454
       PACKAGE_NAME ckprofiler_${GPU_ARCH}
455
456
457
    )
    add_subdirectory(profiler)
  endif()
458
endif()
JD's avatar
JD committed
459
460
461
462
463
464
465
466
467
468

#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
469
configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in
470
471
472
    "${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
473
474
)

475
rocm_install(FILES
JD's avatar
JD committed
476
477
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake"
Anthony Chang's avatar
Anthony Chang committed
478
    DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
JD's avatar
JD committed
479
)
480

481
# Install CK version and configuration files
482
rocm_install(FILES
483
484
485
486
487
    ${PROJECT_BINARY_DIR}/include/ck/version.h
    ${PROJECT_BINARY_DIR}/include/ck/config.h
    DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/ck/
)

488
489
490
491
492
493
494
495
496
497
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
)