[BUG] `#include "cutlass/gemm/device/gemm_universal_adapter.h"` is causing the named symbol to not be found

[mrakgr]

Describe the bug

In the following example, including the #include "cutlass/gemm/device/gemm_universal_adapter.h" line is causing CuPy to be unable to find the qwert_entry0 function. Could including that header be affecting the function names in the compiled program?

Steps/Code to reproduce bug

kernel = r"""
#include <new>
#include <assert.h>
#include <stdio.h>
#include <cutlass/cutlass.h>
#include <cutlass/numeric_types.h>
#include "cutlass/gemm/device/gemm_universal_adapter.h"
using namespace cute;

#include <cooperative_groups.h>
#include <cuda/semaphore>
__device__ cuda::binary_semaphore<cuda::thread_scope_system> console_lock(1);

extern "C" __global__ void qwert_entry0() {
    int v0;
    v0 = threadIdx.x;
    int v1;
    v1 = blockIdx.x;
    int v2;
    v2 = v1 * 256l;
    int v3;
    v3 = v0 + v2;
    bool v4;
    v4 = v3 == 0l;
    if (v4){
        cuda::counting_semaphore<cuda::thread_scope_system, 1l> & v5 = console_lock;
        auto v6 = cooperative_groups::coalesced_threads();
        v5.acquire();
        printf("%s\n","hello");
        v5.release();
        v6.sync() ;
        return ;
    } else {
        return ;
    }
}
"""

import cupy as cp
from dataclasses import dataclass
from typing import NamedTuple, Union, Callable, Tuple
i8 = int; i16 = int; i32 = int; i64 = int; u8 = int; u16 = int; u32 = int; u64 = int; f32 = float; f64 = float; char = str; string = str

options = []
options.append('--define-macro=NDEBUG')
options.append('--dopt=on')
options.append('--diag-suppress=550,20012,68,39,177')
options.append('--restrict')
options.append('--maxrregcount=256')
options.append('-I"G:/cutlass-3.5.1/include"')
options.append('-I"G:/cutlass-3.5.1/tools/util/include"')
options.append('--std=c++20')
options.append('-D__CUDA_NO_HALF_CONVERSIONS__')
raw_module = cp.RawModule(code=kernel, backend='nvcc', enable_cooperative_groups=True, options=tuple(options))
def main_body():
    v0 = cp.cuda.Device().attributes['MultiProcessorCount']
    v1 = v0 == 24
    del v0
    v2 = v1 == False
    if v2:
        v3 = "The number of SMs per GPU at runtime must much that what is declared atop of corecuda.base. Make sure to use the correct constant so it can be propagated at compile time."
        assert v1, v3
        del v3
    else:
        pass
    del v1, v2
    v5 = raw_module.get_function(f"qwert_entry0")
    v5.max_dynamic_shared_size_bytes = 81920 
    v5((24,),(256,),(),shared_mem=81920)
    del v5
    return 

def main():
    r = main_body()
    cp.cuda.get_current_stream().synchronize() # This line is here so the `__trap()` calls on the kernel aren't missed.
    return r

if __name__ == '__main__': print(main())

Expected behavior
Thread 0 should print hello.

Environment details (please complete the following information):

• Environment location: Local, Win 11

Additional context

Here is what happens when I run the script.

PS C:\Spiral_s_ML_Library>  c:; cd 'c:\Spiral_s_ML_Library'; & 'c:\Users\mrakg\AppData\Local\pypoetry\Cache\virtualenvs\ui-EoO7T__V-py3.11\Scripts\python.exe' 'c:\Users\mrakg\.vscode\extensions\ms-python.debugpy-2024.10.0-win32-x64\bundled\libs\debugpy\adapter/../..\debugpy\launcher' '60529' '--' 'c:\Spiral_s_ML_Library\tests\cutlass\test2.py' 
Traceback (most recent call last):
  File "c:\Spiral_s_ML_Library\tests\cutlass\test2.py", line 78, in <module>
    if __name__ == '__main__': print(main())
                                     ^^^^^^
  File "c:\Spiral_s_ML_Library\tests\cutlass\test2.py", line 74, in main
    r = main_body()
  File "c:\Spiral_s_ML_Library\tests\cutlass\test2.py", line 67, in main_body
    v5 = raw_module.get_function(f"qwert_entry0")
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "cupy\_core\raw.pyx", line 487, in cupy._core.raw.RawModule.get_function
  File "cupy\_core\raw.pyx", line 100, in cupy._core.raw.RawKernel.kernel.__get__
  File "cupy\_core\raw.pyx", line 121, in cupy._core.raw.RawKernel._kernel
  File "cupy\cuda\function.pyx", line 275, in cupy.cuda.function.Module.get_function
  File "cupy\cuda\function.pyx", line 216, in cupy.cuda.function.Function.__init__
  File "cupy_backends\cuda\api\driver.pyx", line 244, in cupy_backends.cuda.api.driver.moduleGetFunction
  File "cupy_backends\cuda\api\driver.pyx", line 63, in cupy_backends.cuda.api.driver.check_status
cupy_backends.cuda.api.driver.CUDADriverError: CUDA_ERROR_NOT_FOUND: named symbol not found

In order to actually run the script, you'll have to install CuPy.

[mrakgr]

If you want to try installing CuPy use the pip install cupy-cuda12x command. Trying to install just cupy won't work.

[asi1024]

Here is a minimal reproducer that does not depend on CuPy:

• source

#include "cutlass/gemm/device/gemm_universal_adapter.h"

extern "C" __global__ void my_func() {}

• commands

$ nvcc -I"[CUTLASS_PATH]/include" -I"[CUTLASS_PATH]/tools/util/include" --std=c++17 --device-c -o sample.o sample.cu
$ nvcc -I"[CUTLASS_PATH]/include" -I"[CUTLASS_PATH]/tools/util/include" --cubin --device-link sample.o -o sample.cubin

$ cuobjdump -symbols sample.o

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit
compressed

symbols:
STT_CUDA_OBJECT  STB_LOCAL  STO_GLOBAL     __nv_static_38__ef1904b4_9_sample_cu_b232a47d_2970605__ZN47_INTERNAL_ef1904b4_9_sample_cu_b232a47d_29706054cute1_E
STT_CUDA_OBJECT  STB_LOCAL  STO_GLOBAL     __nv_static_38__ef1904b4_9_sample_cu_b232a47d_2970605__ZN47_INTERNAL_ef1904b4_9_sample_cu_b232a47d_29706054cute7productE
STT_CUDA_OBJECT  STB_LOCAL  STO_?          _SREG
STT_FUNC         STB_GLOBAL STO_ENTRY      my_func

Fatbin ptx code:
================
arch = sm_52
code version = [8,4]
host = linux
compile_size = 64bit
compressed
ptxasOptions = --compile-only

$ cuobjdump -symbols sample.cubin

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

symbols:
STT_OBJECT       STB_LOCAL  STV_DEFAULT    __nv_static_38__ef1904b4_9_sample_cu_b232a47d_2970605__ZN47_INTERNAL_ef1904b4_9_sample_cu_b232a47d_29706054cute1_E
STT_OBJECT       STB_LOCAL  STV_DEFAULT    __nv_static_38__ef1904b4_9_sample_cu_b232a47d_2970605__ZN47_INTERNAL_ef1904b4_9_sample_cu_b232a47d_29706054cute7productE

The file 'sample.o' contains a symbol named my_func, but 'sample.cubin' does not.

[mrakgr]

I'll go ahead and also open an issue with Nvidia for this. I am not a C++ expert, but it's hard for me to imagine that the library itself is doing something to drop the externs. Most likely, this is a NVCC compiler bug.

[mrakgr]

https://developer.nvidia.com/bugs/4862676

[leofang]

Thanks, @mrakgr! One more question, I assume you've tried with NVRTC and you hit the same error, thus switching to NVCC?

[mrakgr]

No, I haven't tried NVRTC. The trouble with NVRTC is that it cannot compile recursive types properly.

https://developer.nvidia.com/bugs/4704632
https://github.com/mrakgr/Spiral-s-ML-Library/blob/9e030d00d50ca9fe6ddcd9bcb39cce0dab2b9b81/tests/test2.py#L182

This example wouldn't compile with NVRTC, but it does with NVCC and as far as I can tell, it's impossible to define recursive union types in NVRTC, so since then I've been using NVCC. The Nvidia rep said they'd fix it.

[mrakgr]

Got a reply from the Cuda rep. For those who cannot see it.

Hi Marko ,
This is expected behavior due to my_func not being referenced anywhere, so it is removed as an optimization. If you add a call to my_func it will remain. Alternatively you could pass -Xnvlink -ignore-host-info to the -dlink command, and that will preserve my_func. See https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#ignore-host-info-ignore-host-info

4.2.9.2.13. --ignore-host-info (-ignore-host-info)
Ignore information about host references, so don’t remove device code that could potentially be referenced by host.

This optimization is useful for some large libraries where there are lots of kernels but only one is used in the application, in this way we save the space of processing all the unused kernels. Note that this optimization of removing unused kernels was introduced in 11.x timeframe, before that we only removed unused device functions (not kernels).
Best,
Yuki

In addition to what he said, maybe --kernels-used would also be good for preserving the function. At any rate, I am not getting aid from the Cutlass authors on how to use the library directly from device code, so I am not going to be testing this or using the library until the situation changes.

[d-k-b]

Hi @mrakgr, I don't see why including that header would affect the visibility of the kernel entry, as you said most likely it is just causing some other heuristic in the compiler to trip and changing the linkage behavior. As mentioned by the CUDA team, it is good practice to reference any kernel you might use from some reachable code in your library to prevent the compiler from removing them.