OpenACC - CUDA interop with CMake

Accelerated Computing HPC Compilers nvc, nvc++ and nvfortran
1

Hi!

We would like to use C with OpenACC pragmas and CUDA kernels in some hotspots.
With make everything is fine. However we would like to build with CMake.

We started with this repo:
https://github.com/jefflarkin/openacc-interoperability

So far we came up with this CMakeLists.txt file:

bash:
export CC=nvc
export CXX=nvc++
export CUDACXX=nvcc

cmake:
cmake_minimum_required(VERSION 3.26)
project(openacc-interoperability)
set(CMAKE_CUDA_ARCHITECTURES 80 86)
set(CMAKE_CUDA_LINK_EXECUTABLE "<CMAKE_CXX_COMPILER> <LINK_FLAGS> <OBJECTS> -o <TARGET>")
enable_language(CUDA)
set(NVC_GPU_LIST ${CMAKE_CUDA_ARCHITECTURES})
list(TRANSFORM NVC_GPU_LIST PREPEND "cc")
list(JOIN NVC_GPU_LIST "," NVC_GPU_LIST)
add_compile_options($<$<COMPILE_LANGUAGE:C,CXX>:-Minfo=accel>)
add_compile_options($<$<COMPILE_LANGUAGE:C,CXX>:-acc>)
add_compile_options($<$<COMPILE_LANGUAGE:C,CXX>:-gpu=${NVC_GPU_LIST}>)
add_link_options(-Minfo=accel -acc -cuda -gpu=${NVC_GPU_LIST})
add_executable(openacc_c_main
    saxpy_cuda.cu
    openacc_c_main.c
)

We needed to remove <LINK_LIBRARIES> from CMAKE_CUDA_LINK_EXECUTABLE because it tries to link cuda libraries in the nvc++ link step, causing errors at runtime. This way however we can not link external libraries, which we want to do.

Our other approach was to compile the CUDA code as a library, and then link it all together. This did not work eighter. The expected to work elegant solution:

cmake_minimum_required(VERSION 3.25)
project(openacc_cuda_interop)
enable_language(CUDA)
find_package(OpenACC REQUIRED)
add_library(saxpy_cuda saxpy_cuda.cu)
add_executable(openacc_c_main openacc_c_main.c)
target_link_libraries(openacc_c_main PRIVATE OpenACC::OpenACC_C)
target_link_libraries(openacc_c_main PRIVATE saxpy_cuda)

What is the correct way of using C with OpenACC, with some CUDA kernels with CMake?

2

Hi Janos,

Unfortunately I wont be able to help with CMake questions since that’s not my area of expertise. Though hopefully can help with more general questions.

We needed to remove <LINK_LIBRARIES> from CMAKE_CUDA_LINK_EXECUTABLE because it tries to link cuda libraries in the nvc++ link step, causing errors at runtime.

I’m not clear on this statement. If you want to link a CUDA object into the binary, then you do need to include the CUDA runtime libraries. The “-cuda” convenience flag does this for you with the added benefit of using the CUDA version matching the one used for compilation. If you didn’t include the CUDA libraries, the executable shouldn’t have linked.

I just tried manually building the binary and don’t see a runtime error. Can you give details on the error you’re seeing?

% nvcc -gencode arch=compute_90,code=sm_90 -c saxpy_cuda.cu
% nvc -cuda -acc -w -c++libs -Minfo=accel saxpy_cuda.o openacc_c_main.c
openacc_c_main.c:
main:
     16, Generating create(x[:n]) [if not already present]
         Generating copyout(y[:n]) [if not already present]
     19, Loop is parallelizable
         Generating NVIDIA GPU code
         19, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
% a.out
y[0] = 2.000000

Our other approach was to compile the CUDA code as a library, and then link it all together. This did not work eighter.

Can you give more details?

As you creating static (.a) or dynamic (.so) libraries?
What didn’t work? Build error? Link error? Runtime error?

-Mat

3

Which cmakefile have you used?
If the first:
Try adding a lib to an executable, run cmake in verbose mode and notice how it does not get linked.

Thank you for the speedy answer!

4

CMakeLists.txt (772 Bytes)
Makefile.txt (1.0 KB)
openacc_c_main.c.txt (1002 Bytes)
saxpy_cuda.cu.txt (1.0 KB)

Uploaded a minimal example with a failing link step. All files have .txt added so I can upload them.
With the makefile you can add -lz to the linker LDFLAGS and it will link correctly.
With CMake either the cuda linking fails with <LINK_LIBRARIES> there (acc_deviceptr returning host pointer, causing cudaIllegalAddress) or zlib linking fails with removed <LINK_LIBRARIES>.

The CMake line with added LINK_LIBRARIES:

set(CMAKE_CUDA_LINK_EXECUTABLE "<CMAKE_CXX_COMPILER> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")

It would be preferable to have an elegant solution work, without the cmake list mangling for example. Is it possible?

Thank you for your speedy answer again, and looking forward to resolve this issue.

5

The error is because you’re missing the zlib library on the link line. To fix, add “-lz” to your link options or undefine “USE_ZLIB” in the openacc_c_main.c file. See your Makefile where you do use “-lz”.

6 
find_package(ZLIB)
target_link_libraries(openacc_c_main ZLIB::ZLIB)

This part should add -lz. It can not, because <LINK_LIBRARIES> is removed from the CMAKE_CUDA_LINK_EXECUTABLE variable. If I add it back then cuda IllegalAddress will occur when the code execution reaches saxpy.

7

Apologies, I misread you’re last post and was just looking at the example. When adding “<LINK_LIBRARIES>”, I can recreate your error.

The issue is that it is adding not only zilb but also several other libraries. Given the order on the link line does matter, by explicitly adding the libraries, they are in the incorrect order. If you manually edit “CMakeFiles/openacc_c_main.dir/linkLibs.rsp” and remove “-lcudadevrt”, “-lcudart_static”, and “-lnvc”, then it will link and run correctly.

On how to fix the cmake LINK_LIBRARIES so the extra libraries aren’t included, or how to just add zlib, you many need to ask the folks at kitware: Code - CMake Discourse

The short-term work around would be to not use LINK_LIBRARIES and add “-lz” to the link options. Though, I don’t know how that would work for your larger project.

8

Well, the end goal is to incorporate hand written cuda kernels in openacc accelerated C/C++ sources and build everything with cmake. We’re not sure if it’s a cmake or a cuda issue, to be fair. My approach was to just write a simple, modern cmake file, my colleague also tried to hack together some of the link lines and so on to try to solve the issues, though ideally that should not be necessary. In any case, building with cmake does not work, but it can be built by hand.

For reference, saxpy_cuda.cu:

__global__
void saxpy_kernel(int n, float a, float *x, float *y)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  if ( i < n )
    y[i] += a * x[i];
}
extern "C" void saxpy(int n ,float a, float *x, float *y)
{
  dim3 griddim, blockdim;

  blockdim = dim3(128,1,1);
  griddim = dim3(n/blockdim.x,1,1);

  saxpy_kernel<<<griddim,blockdim>>>(n,a,x,y);
}

and openacc_c_main.c:

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>

extern void saxpy(int,float,float*,float*);

int main(int argc, char **argv)
{
  float *x, *y;
  int n = 1<<20, i;

  x = (float*)malloc(n*sizeof(float));
  y = (float*)malloc(n*sizeof(float));

  #pragma acc data create(x[0:n]) copyout(y[0:n])
  {
    #pragma acc kernels
    {
      for( i = 0; i < n; i++)
      {
        x[i] = 1.0f;
        y[i] = 0.0f;
      }
    }

    #pragma acc host_data use_device(x,y)
    {
      saxpy(n, 2.0, x, y);
    }
  }

  fprintf(stdout, "y[0] = %f\n",y[0]);
  return 0;
}

To compile and run on Ubuntu 20 and the HPC SDK 23.9 and it’s embedded nvcc

nvcc -gencode arch=compute_80,code=sm_80 -c saxpy_cuda.cu
nvc -fast -acc -gpu=cc80 -c openacc_c_main.c
nvc++ -o openacc_c_main -fast -acc -gpu=cc80 saxpy_cuda.o openacc_c_main.o -cuda

./openacc_c_main
y[0] = 2.000000

But not with cmake and this CMakeLists.txt:

cmake_minimum_required(VERSION 3.25)
project(openacc_cuda_interop)
enable_language(CUDA)

set(CMAKE_EXPORT_COMPILE_COMMANDS true)

find_package(OpenACC REQUIRED)

add_library(saxpy_cuda src/saxpy_cuda.cu)

add_executable(openacc_c_main src/openacc_c_main.c)
target_link_libraries(openacc_c_main PRIVATE OpenACC::OpenACC_C)
target_link_libraries(openacc_c_main PRIVATE saxpy_cuda)

Building with cmake looks like this:

$ CC=nvc CXX=nvc++ cmake -B build . --fresh --warn-uninitialized -DCMAKE_BUILD_TYPE=Debug && cmake --build build/ --config Debug --verbose
Warn about uninitialized values.
-- The C compiler identification is NVHPC 23.9.0
-- The CXX compiler identification is NVHPC 23.9.0
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvc - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvc++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- The CUDA compiler identification is NVIDIA 12.2.91
-- Detecting CUDA compiler ABI info
-- Detecting CUDA compiler ABI info - done
-- Check for working CUDA compiler: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvcc - skipped
-- Detecting CUDA compile features
-- Detecting CUDA compile features - done
-- Found OpenACC_C: -acc
-- Found OpenACC_CXX: -acc
-- Configuring done (5.4s)
-- Generating done (0.0s)
-- Build files have been written to: /home/user/git/openacc-cuda-interop/build
Change Dir: '/home/user/git/openacc-cuda-interop/build'

Run Build Command(s): /usr/bin/cmake -E env VERBOSE=1 /usr/bin/make -f Makefile
/usr/bin/cmake -S/home/user/git/openacc-cuda-interop -B/home/user/git/openacc-cuda-interop/build --check-build-system CMakeFiles/Makefile.cmake 0
/usr/bin/cmake -E cmake_progress_start /home/user/git/openacc-cuda-interop/build/CMakeFiles /home/user/git/openacc-cuda-interop/build//CMakeFiles/progress.marks
/usr/bin/make  -f CMakeFiles/Makefile2 all
make[1]: Entering directory '/home/user/git/openacc-cuda-interop/build'
/usr/bin/make  -f CMakeFiles/saxpy_cuda.dir/build.make CMakeFiles/saxpy_cuda.dir/depend
make[2]: Entering directory '/home/user/git/openacc-cuda-interop/build'
cd /home/user/git/openacc-cuda-interop/build && /usr/bin/cmake -E cmake_depends "Unix Makefiles" /home/user/git/openacc-cuda-interop /home/user/git/openacc-cuda-interop /home/user/git/openacc-cuda-interop/build /home/user/git/openacc-cuda-interop/build /home/user/git/openacc-cuda-interop/build/CMakeFiles/saxpy_cuda.dir/DependInfo.cmake "--color="
make[2]: Leaving directory '/home/user/git/openacc-cuda-interop/build'
/usr/bin/make  -f CMakeFiles/saxpy_cuda.dir/build.make CMakeFiles/saxpy_cuda.dir/build
make[2]: Entering directory '/home/user/git/openacc-cuda-interop/build'
[ 25%] Building CUDA object CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o
/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvcc -forward-unknown-to-host-compiler   -g "--generate-code=arch=compute_52,code=[compute_52,sm_52]" -MD -MT CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o -MF CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o.d -x cu -c /home/user/git/openacc-cuda-interop/src/saxpy_cuda.cu -o CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o
[ 50%] Linking CUDA static library libsaxpy_cuda.a
/usr/bin/cmake -P CMakeFiles/saxpy_cuda.dir/cmake_clean_target.cmake
/usr/bin/cmake -E cmake_link_script CMakeFiles/saxpy_cuda.dir/link.txt --verbose=1
/usr/bin/ar qc libsaxpy_cuda.a CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o
/usr/bin/ranlib libsaxpy_cuda.a
make[2]: Leaving directory '/home/user/git/openacc-cuda-interop/build'
[ 50%] Built target saxpy_cuda
/usr/bin/make  -f CMakeFiles/openacc_c_main.dir/build.make CMakeFiles/openacc_c_main.dir/depend
make[2]: Entering directory '/home/user/git/openacc-cuda-interop/build'
cd /home/user/git/openacc-cuda-interop/build && /usr/bin/cmake -E cmake_depends "Unix Makefiles" /home/user/git/openacc-cuda-interop /home/user/git/openacc-cuda-interop /home/user/git/openacc-cuda-interop/build /home/user/git/openacc-cuda-interop/build /home/user/git/openacc-cuda-interop/build/CMakeFiles/openacc_c_main.dir/DependInfo.cmake "--color="
make[2]: Leaving directory '/home/user/git/openacc-cuda-interop/build'
/usr/bin/make  -f CMakeFiles/openacc_c_main.dir/build.make CMakeFiles/openacc_c_main.dir/build
make[2]: Entering directory '/home/user/git/openacc-cuda-interop/build'
[ 75%] Building C object CMakeFiles/openacc_c_main.dir/src/openacc_c_main.c.o
/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvc   -g -O0 -acc -MD -MT CMakeFiles/openacc_c_main.dir/src/openacc_c_main.c.o -MF CMakeFiles/openacc_c_main.dir/src/openacc_c_main.c.o.d -o CMakeFiles/openacc_c_main.dir/src/openacc_c_main.c.o -c /home/user/git/openacc-cuda-interop/src/openacc_c_main.c
[100%] Linking CUDA executable openacc_c_main
/usr/bin/cmake -E cmake_link_script CMakeFiles/openacc_c_main.dir/link.txt --verbose=1
/usr/bin/g++ @CMakeFiles/openacc_c_main.dir/objects1.rsp -o openacc_c_main @CMakeFiles/openacc_c_main.dir/linkLibs.rsp -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/lib64" -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/math_libs/12.2/lib64" -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/comm_libs/12.2/nccl/lib" -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/comm_libs/12.2/nvshmem/lib" -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/targets/x86_64-linux/lib/stubs" -L"/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/cuda/12.2/targets/x86_64-linux/lib"
/usr/bin/ld: CMakeFiles/openacc_c_main.dir/src/openacc_c_main.c.o: in function `main':
/home/user/git/openacc-cuda-interop/src/openacc_c_main.c:24: undefined reference to `__pgi_uacc_enter'
/usr/bin/ld: /home/user/git/openacc-cuda-interop/src/openacc_c_main.c:18: undefined reference to `__pgi_uacc_computestart2'
/usr/bin/ld: /home/user/git/openacc-cuda-interop/src/openacc_c_main.c:19: undefined reference to `__pgi_uacc_launch'
/usr/bin/ld: /home/user/git/openacc-cuda-interop/src/openacc_c_main.c:24: undefined reference to `__pgi_uacc_computedone'
/usr/bin/ld: /home/user/git/openacc-cuda-interop/src/openacc_c_main.c:24: undefined reference to `__pgi_uacc_noversion'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_cooperlake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_get_l3_cachesize'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_rocketlake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_remainder'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_haswell'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_icelake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_j1'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_znver4'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_tigerlake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_avx2'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_y0'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_avx512vl'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_znver3'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_j0'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dremainder'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_get_threads_per_socket'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dround'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_sandybridge'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_y1'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_gh_b'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_znver2'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_y1'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_intel'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_jn'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_sapphirerapids'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_znver1'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_y0'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_jn'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_cascadelake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_skylake_xeon'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_avx512f'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_skylake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_dbessel_yn'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_alderlake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_fma'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_ivybridge'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_graniterapids'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_j0'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_cannonlake'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_j1'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_skylake_client'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_around'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__mth_i_bessel_yn'
/usr/bin/ld: /opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/lib/libnvc.so: undefined reference to `__Cpuid_is_broadwell'
collect2: error: ld returned 1 exit status
make[2]: *** [CMakeFiles/openacc_c_main.dir/build.make:100: openacc_c_main] Error 1
make[2]: Leaving directory '/home/user/git/openacc-cuda-interop/build'
make[1]: *** [CMakeFiles/Makefile2:111: CMakeFiles/openacc_c_main.dir/all] Error 2
make[1]: Leaving directory '/home/user/git/openacc-cuda-interop/build'
make: *** [Makefile:91: all] Error 2

Any idea what is going on and how to fix this on the cmake side? It looks like it’s using g++ and ld for linking, does that make sense? Is cmake maybe confused what the link language should be for linking a cuda library and a C executable?

9

As I noted above, I don’t use cmake much myself and don’t really have great insights.

However, from what I can tell from the previous poster’s CMakeLists.txt, you need to set the “CMAKE_CUDA_LINK_EXECUTABLE” to use the CXX compiler. Looks like it defaults to using g++. I also needed to add some link flags:

cmake_minimum_required(VERSION 3.25)
project(openacc_cuda_interop)
set(CMAKE_EXPORT_COMPILE_COMMANDS true)
enable_language(CUDA)
set(CMAKE_CUDA_LINK_EXECUTABLE "<CMAKE_CXX_COMPILER> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
find_package(OpenACC REQUIRED)
add_library(saxpy_cuda src/saxpy_cuda.cu)
add_link_options(-Minfo=accel -acc -cuda)
add_executable(openacc_c_main src/openacc_c_main.c)
target_link_libraries(openacc_c_main PRIVATE OpenACC::OpenACC_C)
target_link_libraries(openacc_c_main PRIVATE saxpy_cuda)

The generated binary wont error like above but doesn’t seem to launch the CUDA kernel correctly. Looks to be the same problem as above where the extra runtime libraries Cmake is adding to “LINK_LIBRARIES” is causing issue. Remove the extra libraries from “CMakeFiles/openacc_c_main.dir/linkLibs.rsp” and it will get correct results.

There might be a way to manipulate “LINK_LIBRARIES” so the extra runtime libraries aren’t included, but I’m not sure. It seems to me that Cmake assumes it’s linking with g++ which doesn’t implicitly include the dependent CUDA libraries. Again for the OpenACC/CUDA interoperability, the order in which the runtime libraries appear in the link matters, so only nvc++ should be adding them so they are in the correct order.

10

Does this even preserve the device code? If not I suspect that is why no cuda kernels are called and the result is wrong even with your suggestion of setting the link executable and the two additional link options.

/usr/bin/ar qc libsaxpy_cuda.a CMakeFiles/saxpy_cuda.dir/src/saxpy_cuda.cu.o
/usr/bin/ranlib libsaxpy_cuda.a

I can also confirm that setting -cuda and -acc as link options for the executable (the modern cmake way is to not set this globally) get’s rid of the missing symbols, but… -acc is a link option that is inherited from linking against the imported target OpenACC::OpenACC_C and I have confirmed that it is set there. So it should not be necessary as it should already be on the link line. That is a cmake question though.

Maybe the best solution as it both compiles and works correctly / launches kernels: setting the cuda parts to be a shared library. That works even without manually adding -cuda and -acc as link options and without setting the link executable, so basically with my simple example above, except this:

add_library(saxpy_cuda SHARED src/saxpy_cuda.cu)

That is not really what we want though… If that is needed then ideally the cuda and openacc parts both form a shared lib. But we can work with that if needs be.

My grief though with all this is that I do NOT want to deal with all of this. This is why you use cmake with it’s CUDA language support and things like the OpenACC imported target from the according package. I know that it is possible to hack all of this, but that is exactly not what you want to do with modern cmake. If I wanted that I would use make. ;) Cmake has a lot of information and I expect it to use that information to make my life easier. That is (again) a cmake topic however. We will ask it there and come back with follow-up questions potentially.