CMakeLists.txt 18.6 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
12
13
14
15
16
17
18
19
20
if(_GENERATOR_IS_MULTI_CONFIG)
    if (NOT CMAKE_CONFIGURATION_TYPES)
        set(CMAKE_CONFIGURATION_TYPES "Debug;Release;RelWithDebInfo;MinSizeRel" CACHE STRING
                "Available build types (configurations) on multi-config generators")
    endif()
else()
    if(NOT CMAKE_BUILD_TYPE)
        set(CMAKE_BUILD_TYPE Release CACHE STRING
                "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel.")
    endif()
21
22
23
endif()

# Default installation path
Artur Wojcik's avatar
Artur Wojcik committed
24
if(NOT WIN32)
25
26
27
    set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "")
endif()

28
set(version 1.1.0)
JD's avatar
JD committed
29
# Check support for CUDA/HIP in Cmake
Artur Wojcik's avatar
Artur Wojcik committed
30
31
project(composable_kernel VERSION ${version} LANGUAGES CXX)
include(CTest)
Chao Liu's avatar
Chao Liu committed
32

33
34
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")

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

71
72
73
#for f8/bf8_t type
add_compile_options(-Wno-bit-int-extension)

74
75
if(DL_KERNELS)
    add_definitions(-DDL_KERNELS)
76
    set(CK_ENABLE_DL_KERNELS "ON")
77
78
79
80
endif()

if(INSTANCES_ONLY)
    add_definitions(-DINSTANCES_ONLY)
81
    set(CK_ENABLE_INSTANCES_ONLY "ON")
82
83
endif()

Artur Wojcik's avatar
Artur Wojcik committed
84
85
include(GetOpt)

86
# CK config file to record supported datatypes, etc.
Artur Wojcik's avatar
Artur Wojcik committed
87
configure_file(include/ck/config.h.in ${CMAKE_CURRENT_BINARY_DIR}/include/ck/config.h)
88
89
90
91

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

94
set(ROCM_SYMLINK_LIBS OFF)
Anthony Chang's avatar
Anthony Chang committed
95
find_package(ROCM REQUIRED PATHS /opt/rocm)
JD's avatar
JD committed
96
97
98
99
100
101

include(ROCMInstallTargets)
include(ROCMPackageConfigHelpers)
include(ROCMSetupVersion)
include(ROCMInstallSymlinks)
include(ROCMCreatePackage)
Chao Liu's avatar
Chao Liu committed
102
include(CheckCXXCompilerFlag)
103
include(ROCMCheckTargetIds)
JD's avatar
JD committed
104
include(TargetFlags)
105
106
107

rocm_setup_version(VERSION ${version})

Artur Wojcik's avatar
Artur Wojcik committed
108
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
109

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

139
message("Supported GPU_TARGETS= ${DEFAULT_GPU_TARGETS}")
140

141
set(AMDGPU_TARGETS "${DEFAULT_GPU_TARGETS}" CACHE STRING " " FORCE)
142
143
144
145
146
147

if(GPU_TARGETS)
    message("Building CK for the following targets: ${GPU_TARGETS}")
else()
    message("Building CK for the following targets: ${AMDGPU_TARGETS}")
endif()
148
## C++
Chao Liu's avatar
Chao Liu committed
149
set(CMAKE_CXX_STANDARD 17)
Chao Liu's avatar
Chao Liu committed
150
set(CMAKE_CXX_STANDARD_REQUIRED ON)
Chao Liu's avatar
Chao Liu committed
151
set(CMAKE_CXX_EXTENSIONS OFF)
Chao Liu's avatar
Chao Liu committed
152
153
message("CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")

154
option(CK_BUILD_JIT_LIB "Only build the CK JIT Helper Library" OFF)
Alan Turner's avatar
Alan Turner committed
155
if (NOT CK_BUILD_JIT_LIB)
156
157
158
159
160
161
162
163
164
    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}")
    if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 500723302)
       message("Adding the fno-offload-uniform-block compiler flag")
       add_compile_options(-fno-offload-uniform-block)
    endif()
Chao Liu's avatar
Chao Liu committed
165

Alan Turner's avatar
Alan Turner committed
166
    option(USE_BITINT_EXTENSION_INT4, "Whether to enable clang's BitInt extension to provide int4 data type." OFF)
Alan Turner's avatar
Alan Turner committed
167
    option(USE_OPT_NAVI3X, "Whether to enable LDS cumode and Wavefront32 mode for NAVI3X silicons." OFF)
Chao Liu's avatar
Chao Liu committed
168

Alan Turner's avatar
Alan Turner committed
169
170
171
172
173
    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()
Chao Liu's avatar
Chao Liu committed
174

Alan Turner's avatar
Alan Turner committed
175
176
177
178
179
    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()
Chao Liu's avatar
Chao Liu committed
180

Alan Turner's avatar
Alan Turner committed
181
    ## Threads
182
183
184
    if(NOT WIN32)
        set(THREADS_PREFER_PTHREAD_FLAG ON)
    endif()
Alan Turner's avatar
Alan Turner committed
185
186
    find_package(Threads REQUIRED)
    link_libraries(Threads::Threads)
187

Alan Turner's avatar
Alan Turner committed
188
189
    ## OpenMP
    if(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
190
191
192
193
194
195
196
	    # 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})
Alan Turner's avatar
Alan Turner committed
197
    else()
198
	    find_package(OpenMP REQUIRED)
Alan Turner's avatar
Alan Turner committed
199
    endif()
Chao Liu's avatar
Chao Liu committed
200

Alan Turner's avatar
Alan Turner committed
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
    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}")

    link_libraries(${OpenMP_gomp_LIBRARY})
    link_libraries(${OpenMP_pthread_LIBRARY})

    ## HIP
    find_package(HIP REQUIRED)
    # 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}")
    link_libraries(hip::device)
    add_compile_definitions(__HIP_PLATFORM_HCC__=1)
JD's avatar
JD committed
232
endif()
233

Chao Liu's avatar
Chao Liu committed
234
235
## tidy
include(EnableCompilerWarnings)
JD's avatar
JD committed
236
set(CK_TIDY_ERRORS ERRORS * -readability-inconsistent-declaration-parameter-name)
Chao Liu's avatar
Chao Liu committed
237
if(CMAKE_CXX_COMPILER MATCHES ".*hcc" OR CMAKE_CXX_COMPILER MATCHES ".*clang\\+\\+")
JD's avatar
JD committed
238
    set(CK_TIDY_CHECKS -modernize-use-override -readability-non-const-parameter)
Chao Liu's avatar
Chao Liu committed
239
# Enable tidy on hip
JD's avatar
JD committed
240
241
elseif(CK_BACKEND STREQUAL "HIP" OR CK_BACKEND STREQUAL "HIPNOGPU")
    set(CK_TIDY_ERRORS ALL)
Chao Liu's avatar
Chao Liu committed
242
243
endif()

JD's avatar
JD committed
244

Chao Liu's avatar
Chao Liu committed
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
327
328
329
330
331
332
333
334
335
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
336
337
        ${CK_TIDY_CHECKS}
        ${CK_TIDY_ERRORS}
Chao Liu's avatar
Chao Liu committed
338
339
340
    HEADER_FILTER
        "\.hpp$"
    EXTRA_ARGS
JD's avatar
JD committed
341
        -DCK_USE_CLANG_TIDY
Chao Liu's avatar
Chao Liu committed
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
)

include(CppCheck)
enable_cppcheck(
    CHECKS
        warning
        style
        performance
        portability
    SUPPRESS
        ConfigurationNotChecked
        constStatement
        duplicateCondition
        noExplicitConstructor
        passedByValue
Chao Liu's avatar
Chao Liu committed
357
        preprocessorErrorDirective
Chao Liu's avatar
Chao Liu committed
358
359
360
361
362
363
364
        shadowVariable
        unusedFunction
        unusedPrivateFunction
        unusedStructMember
        unmatchedSuppression
    FORCE
    SOURCES
Chao Liu's avatar
Chao Liu committed
365
        library/src
Chao Liu's avatar
Chao Liu committed
366
367
368
    INCLUDE
        ${CMAKE_CURRENT_SOURCE_DIR}/include
        ${CMAKE_CURRENT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
369
        ${CMAKE_CURRENT_SOURCE_DIR}/library/include
Chao Liu's avatar
Chao Liu committed
370
371
372
373
    DEFINE
        CPPCHECK=1
        __linux__=1
)
Chao Liu's avatar
Chao Liu committed
374

JD's avatar
JD committed
375
376
377
378
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)

379
# set CK project include directories
Chao Liu's avatar
Chao Liu committed
380
include_directories(BEFORE
381
    ${PROJECT_BINARY_DIR}/include
Chao Liu's avatar
Chao Liu committed
382
383
    ${PROJECT_SOURCE_DIR}/include
    ${PROJECT_SOURCE_DIR}/library/include
384
    ${HIP_INCLUDE_DIRS}
JD's avatar
JD committed
385
)
Chao Liu's avatar
Chao Liu committed
386

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

Alan Turner's avatar
Alan Turner committed
389
if (NOT CK_BUILD_JIT_LIB)
Alan Turner's avatar
Alan Turner committed
390
391
    SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV")

Alan Turner's avatar
Alan Turner committed
392
393
394
395
    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})
396
397
398
399
400
    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)
401
        if(("${cmake_instance}" MATCHES "fp8" OR "${cmake_instance}" MATCHES "_f8") AND DTYPES MATCHES "fp8")
402
403
404
            #message("fp8 instance found!")
            set(add_inst 1)
        endif()
405
406
407
408
409
        if(("${cmake_instance}" MATCHES "bf8" OR "${cmake_instance}" MATCHES "_b8") AND DTYPES MATCHES "bf8")
            #message("bf8 instance found!")
            set(add_inst 1)
        endif()
        if(("${cmake_instance}" MATCHES "fp16" OR "${cmake_instance}" MATCHES "_f16") AND DTYPES MATCHES "fp16")
410
411
412
            #message("fp16 instance found!")
            set(add_inst 1)
        endif()
413
        if(("${cmake_instance}" MATCHES "fp32" OR "${cmake_instance}" MATCHES "_f32") AND DTYPES MATCHES "fp32")
414
415
416
            #message("fp32 instance found!")
            set(add_inst 1)
        endif()
417
        if(("${cmake_instance}" MATCHES "fp64" OR "${cmake_instance}" MATCHES "_f64") AND DTYPES MATCHES "fp64")
418
419
420
            #message("fp64 instance found!")
            set(add_inst 1)
        endif()
421
        if(("${cmake_instance}" MATCHES "bf16" OR "${cmake_instance}" MATCHES "_b16") AND DTYPES MATCHES "bf16")
422
423
424
            #message("bf16 instance found!")
            set(add_inst 1)
        endif()
425
        if(("${cmake_instance}" MATCHES "int8" OR "${cmake_instance}" MATCHES "_i8") AND DTYPES MATCHES "int8")
426
427
428
429
430
431
432
433
434
435
436
            #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()
Alan Turner's avatar
Alan Turner committed
437
    ENDFOREACH()
438

Alan Turner's avatar
Alan Turner committed
439
    add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES}  SOURCES ${INSTANCE_FILES})
440
    add_subdirectory(library)
Anthony Chang's avatar
Anthony Chang committed
441

442
443
    if(NOT DEFINED INSTANCES_ONLY)
    if(NOT DEFINED PROFILER_ONLY)
Alan Turner's avatar
Alan Turner committed
444
445
446
447
    rocm_package_setup_component(tests
            LIBRARY_NAME composablekernel
            PACKAGE_NAME tests # Prevent -static suffix on package name
    )
448

Alan Turner's avatar
Alan Turner committed
449
450
451
452
    rocm_package_setup_component(examples
            LIBRARY_NAME composablekernel
            PACKAGE_NAME examples
    )
453
    add_subdirectory(example)
Artur Wojcik's avatar
Artur Wojcik committed
454
   if(BUILD_TESTING)
455
       add_subdirectory(test)
Artur Wojcik's avatar
Artur Wojcik committed
456
   endif()
457

458
    rocm_package_setup_component(profiler
Alan Turner's avatar
Alan Turner committed
459
            LIBRARY_NAME composablekernel
460
        PACKAGE_NAME ckProfiler
461
462
    )
    add_subdirectory(profiler)
463
464
465
466
    else()
        #When building PROFILER_ONLY, label the package with GPU_ARCH
        rocm_package_setup_component(profiler
        LIBRARY_NAME composablekernel
467
       PACKAGE_NAME ckProfiler_${GPU_ARCH}
468
469
470
471
        )
        add_subdirectory(profiler)
    endif()
    endif()
472
else()
Alan Turner's avatar
Alan Turner committed
473
    rocm_package_setup_component(jit_library
474
475
        LIBRARY_NAME composablekernel
        PACKAGE_NAME jit_library
Alan Turner's avatar
Alan Turner committed
476
    )
477
478
    add_subdirectory(library)
    add_subdirectory(test)
479
endif()
JD's avatar
JD committed
480
481
482
483
484
485
486
487
488
489

#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
490
configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in
491
492
493
    "${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
494
495
)

496
rocm_install(FILES
JD's avatar
JD committed
497
498
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake"
    "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake"
Anthony Chang's avatar
Anthony Chang committed
499
    DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel
JD's avatar
JD committed
500
)
501

502
# Install CK version and configuration files
503
rocm_install(FILES
504
505
506
507
508
    ${PROJECT_BINARY_DIR}/include/ck/version.h
    ${PROJECT_BINARY_DIR}/include/ck/config.h
    DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/ck/
)

509
510
511
512
513
514
515
516
517
518
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
)