[BUG] Cutlass' static library defines fewer symbols than the shared library

[joerowell]

Describe the bug
I've recently been playing around with building cutlass as both a shared and static library. However, I've noticed that the static library often omits some symbols that are present in the shared library.

Steps/Code to reproduce bug
I've encountered this issue on 3a8c01a. I've not tried any other cutlass versions, but I don't think there would be a difference.

First, make and navigate to a build directory from the cutlass root directory. Then, configure cutlass' build system using the following cmake command:

cmake .. -DCUTLASS_NVCC_ARCHS=90 -DCUTLASS_ENABLE_TESTS=OFF -DCUTLASS_UNITY_BUILD_ENABLED=ON -DCUTLASS_LIBRARY_OPERATIONS=gemm

The NVCC archs, tests and unity builds aren't that important for reproducibility: I've seen the issue occur in many different configurations of flags.

Then, compile both the static and the shared libraries via make.

make cutlass_library
make cutlass_library_static

This should succeed. Finally, run the following commands and notice the different outputs:

readelf -Ws tools/library/libcutlass.a --demangle | grep cutlass::library::from_string | wc -l ; I get 12
readelf -Ws tools/library/libcutlass.so --demangle | grep cutlass::library::from_string | wc -l ; I get 24

Moreover, if I actually look at the symbols themselves, I see that the static library doesn't actually define all of the symbols.

readelf -Ws tools/library/libcutlass.a --demangle | grep "cutlass::library::initialize_all(cutlass::library::Manifest&)

   9: 0000000000000000     0 NOTYPE  GLOBAL DEFAULT  UND cutlass::library::initialize_all(cutlass::library::Manifest&)
   7: 0000000000000000    27 FUNC    GLOBAL DEFAULT    1 cutlass::library::initialize_all(cutlass::library::Manifest&)

Whereas the shared library does define these:

readelf -Ws tools/library/libcutlass.so --demangle | grep "cutlass::library::intialize_all(cutlass::library::Manifest&)

 13441: 00000000005c5ee0    27 FUNC    GLOBAL DEFAULT   14 cutlass::library::initialize_all(cutlass::library::Manifest&)
 15449: 00000000005c5ee0    27 FUNC    GLOBAL DEFAULT   14 cutlass::library::initialize_all(cutlass::library::Manifest&)

This means that a third-party application that wants to statically link against cutlass is somewhat limited.

Expected behavior
My expectation would be that the static library would export all non-explicitly hidden symbols. I appreciate there's some subtle linker behaviour that can influence how symbols are exported, but this seems unusual.

[joerowell]

We debugged this issue and found out that the static library doesn't appear to be link against the compiled versions of the autogenerated kernels i.e. the files in build/tools/library/ aren't explicitly linked against libcutlass.a.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[d-k-b]

@joerowell, aside from the symbol visibility, if you link against the static libraries does your application function properly? Before spending too much time debugging, I was just curious if this is mostly an academic exercise or if you have noticed failures trying to link against the static versions.

The visibility rules and how the linker handles the libraries is different between static and shared, so sometimes it's not obvious where symbols are found. For instance, a shared library will show symbols that are imported from another library. Libraries don't really link against each other, they just define symbols they import/export and then executables (at build time) and the dynamic linker (at runtime) determine how to patch everything up.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.