[SOLVED] CMake and Cuda separate compilation of class constructor on device fail during linking

Issue

I have problem with linking constructor that is on device via CMake. It causes troubles only when calling function within class. Do you have any idea?

I tried to compile and link manually with nvcc ... -dc and -dlink and that works fine.

Failed output

ptxas fatal   : Unresolved extern function '_ZN4TestC1Ev'
CMake Error at genetic_program_lib_generated_test.cu.o.debug.cmake:282 (message):
  Error generating file
  /article/src/build/CMakeFiles/genetic_program_lib.dir/src/./genetic_program_lib_generated_test.cu.o

CMake

cmake_minimum_required(VERSION 3.22)
project(GeneticProgramming)

set(CMAKE_PREFIX_PATH ${CMAKE_PREFIX_PATH} "/article/src/build")
set(CMAKE_BUILD_TYPE debug)

include_directories(include)

# add cuda source files to variable
file(GLOB CUDA_SOURCES src/genetic_program.cu src/test.cu src/test2.cu)
file(GLOB MAIN main/main.cpp)

# find the cuda package
find_package(CUDA REQUIRED)

cuda_add_library(genetic_program_lib ${CUDA_SOURCES})
set_target_properties(genetic_program_lib PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

# create the executable
add_executable(out ${MAIN})
target_link_libraries(out genetic_program_lib)

Main

#include <genetic_program.hpp>
int main() 
{
    genetic_program();
    return 0;
}

genetic_program.hpp

#pragma once
#include <iostream>
int genetic_program();

genetic_program.cu

#include <genetic_program.hpp>
#include <test.cuh>
int genetic_program()
{
    test<<<1,1>>>();
    return 0;
}

test.cuh

#pragma once
#include <iostream>
#include <test2.cuh>

__global__ void test();

test.cu

#include <test.cuh>
__global__ void test(){
    printf("Hello World!\n");
    Test t;
    t.test2t();
}

test2.cuh

#pragma once
#include <iostream>

class Test
{
    public:
        __device__ Test();
        __device__ void test2t();
};

test2.cu

#include <test2.cuh>
__device__ Test::Test()
{
    printf("Test::Test()\n");
}
__device__ void Test::test2t()
{
    printf("Test::test2t()\n");
}

EDIT: Update naming convention.

Solution

Both CUDA language support and the CUDA_SEPARABLE_COMPILATION property came with CMake 3.8 and don't work together with the legacy CUDA package.

cmake_minimum_required(VERSION 3.22)

set(CMAKE_PREFIX_PATH ${CMAKE_PREFIX_PATH} "/article/src/build")
set(CMAKE_BUILD_TYPE debug)

project(GeneticProgramming LANGUAGES CUDA CXX)

include_directories(include)

# add cuda source files to variable
file(GLOB CUDA_SOURCES src/genetic_program.cu src/test.cu src/test2.cu)
file(GLOB MAIN main/main.cpp)

# cuda_add_library(genetic_program_lib ${CUDA_SOURCES})
add_library(genetic_program_lib ${CUDA_SOURCES})
set_target_properties(genetic_program_lib PROPERTIES CUDA_SEPARABLE_COMPILATION ON) 

# create the executable
add_executable(out ${MAIN})
target_link_libraries(out genetic_program_lib)

works for me.

Note that global variables like CMAKE_BUILD_TYPE need to be set before the project call to take effect.

Answered By - paleonix

Answer Checked By - Candace Johnson (WPSolving Volunteer)