[Issue]: Cannot compile composable_kernel code without -O2 optimization (inline assembly 'n' constraint requires immediate)

[imyixinw]

Problem Description

During compilation with -O2, the following error is observed multiple times in different files: error: constraint 'n' expects an integer constant expression.

Operating System

CentOS Stream 9

CPU

AMD EPYC 9654 96-Core Processor

GPU

AMD Instinct MI300X

ROCm Version

ROCm 6.4.1

ROCm Component

Composable Kernel

Steps to Reproduce

Suppose composable_kernel is already built following the instructions.

Create a new folder (e.g., repro/) containing CMakeLists.txt (download CMakeLists.txt) and minimal_test.hip (download minimal_test.hip). The minimal test is an example of a kernel generated when building the FMHA kernel in PyTorch, which includes the CK library as a dependency. We cannot get a debuggable O0 build of PyTorch without this being fixed.

Under the subdirectory inside the repro folder (e.g., repro/build/), run cmake .. and then make -j.

We should expect to see this error:

$ make -j
[ 50%] Building HIP object CMakeFiles/repro_test.dir/minimal_test.hip.o
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:11:
.../composable_kernel/repro/../include/ck_tile/core/arch/amd_buffer_addressing.hpp:754:18: error: constraint 'n' expects an integer constant expression
  754 |     asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory");
      |                  ^
.../composable_kernel/repro/../include/ck_tile/core/arch/amd_buffer_addressing.hpp:1302:39: error: constraint 'n' expects an integer constant expression
 1302 |         CK_TILE_ASYNC_LOAD_WITH_INSTR("buffer_load_dword");
      |                                       ^
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:16:
.../composable_kernel/repro/../include/ck_tile/core/arch/utility.hpp:27:18: error: constraint 'n' expects an integer constant expression
   27 |     asm volatile("s_add_u32 m0, %0, m0" : : "n"(v) : "memory");
      |                  ^
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:11:
.../composable_kernel/repro/../include/ck_tile/core/arch/amd_buffer_addressing.hpp:996:18: error: constraint 'n' expects an integer constant expression
  996 |     asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory");
      |                  ^
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:45:
.../composable_kernel/repro/../include/ck_tile/core/tensor/load_tile.hpp:124:18: error: constraint 'n' expects an integer constant expression
  124 |     asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory");
      |                  ^
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:11:
.../composable_kernel/repro/../include/ck_tile/core/arch/amd_buffer_addressing.hpp:1002:18: error: constraint 'n' expects an integer constant expression
 1002 |     asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory");
      |                  ^
In file included from .../composable_kernel/repro/minimal_test.hip:8:
In file included from .../composable_kernel/repro/../example/ck_tile/01_fmha/fmha_fwd.hpp:6:
In file included from .../composable_kernel/repro/../include/ck_tile/core.hpp:16:
.../composable_kernel/repro/../include/ck_tile/core/arch/utility.hpp:21:18: error: invalid operand for instruction
   21 |     asm volatile("s_mov_b32 m0, %0" : : "s"(v) : "memory");
      |                  ^
<inline asm>:1:16: note: instantiated into assembly here
    1 |         s_mov_b32 m0, v0
      |                       ^
7 errors generated when compiling for gfx942.
make[2]: *** [CMakeFiles/repro_test.dir/build.make:75: CMakeFiles/repro_test.dir/minimal_test.hip.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:83: CMakeFiles/repro_test.dir/all] Error 2
make: *** [Makefile:91: all] Error 2

The error: constraint 'n' expects an integer constant expression indicates that the inline assembly expects a compile-time immediate value. However, in this case, it uses a value that is a function argument, which is not an immediate. As a result, the code relies on compiler optimizations (e.g. inlining) to make it legal, so it's expectedly failing at -O0. This behavior suggests the issue is an application/library bug, rather than a compilation issue.

[asleepzzz]

we will use constexpr to fix it, and "invalid operand for instruction: is another isssue, you can follow the instruction format

[dmpots]

we will use constexpr to fix it,

Great, thanks for the fix!

"invalid operand for instruction: is another isssue, you can follow the instruction format

Sorry, I don't understand. Do you mean there is a separate issue here that needs a different fix and will not be solved by the constexpr change?