From f8ec330b69f592e764eb4064df823149de223b13 Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Fri, 14 Nov 2025 17:39:50 -0800 Subject: [PATCH] Disable DL kernels on all architectures except gfx103x. (#3218) * disable dl kernels on all archs except gfx103 * add gfx10-3-generic target to cmake [ROCm/composable_kernel commit: b38bb492a1a55b5abb0c345962143c0f9c482cfb] --- CMakeLists.txt | 2 +- include/ck/utility/amd_inline_asm.hpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7b4990dba4..b1532f2cc8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -122,7 +122,7 @@ add_compile_options( # Recent change in compiler makes this warning ON by default, which led to compile errors. add_compile_options(-Wno-nrvo) -if(NOT DISABLE_DL_KERNELS) +if(NOT DISABLE_DL_KERNELS AND GPU_TARGETS MATCHES "gfx103|gfx10-3-generic") add_definitions(-DDL_KERNELS) set(DL_KERNELS "ON") set(CK_ENABLE_DL_KERNELS "ON") diff --git a/include/ck/utility/amd_inline_asm.hpp b/include/ck/utility/amd_inline_asm.hpp index 0ed60df2c3..efe1f300c2 100644 --- a/include/ck/utility/amd_inline_asm.hpp +++ b/include/ck/utility/amd_inline_asm.hpp @@ -86,7 +86,7 @@ inline __device__ f8x8_t amd_assembly_i4_to_fp8x8(int a) return bit_cast(((static_cast(fp8x4_1) << 32) | fp8x4_0)); } - +#ifdef DL_KERNELS // c0 += inner_product(a, b0) // c1 += inner_product(a, b1) __device__ void amd_assembly_outer_product_1x2(float a, float b0, float b1, float& c0, float& c1) @@ -430,6 +430,6 @@ __device__ void amd_assembly_outer_product_1x4(int8x16_t a, c2, c3); } - +#endif } // namespace ck #endif