Cuda C++ template undefined reference Compilation issue while porting a visual studio project to linux

Question

I am trying to compile the following code ([cuShiftOr]) to a linux box. Original code is running on Visual Studio 2013 in Windows.

My strategy is to compile all .cu files like below:

nvcc -std=c++11 -Icpp11-range -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env --cl-version 2013 -maxrregcount=0 --machine 64 --compile -cudart static -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -I$HOME/cuShiftOr/CuShiftOr -I$HOME/cuShiftOr/CuShiftOrBenchmark -I$HOME/cuShiftOr/tinyformat -Xcompiler "-fPIC -fexceptions -ffunction-sections -fdata-sections -fpermissive" -ccbin=/opt/cray/pe/craype/2.5.10/bin/CC -c CuShiftOr/device/hybrid.cu -o objs/CuShiftOr/device/hybrid.cu.o

Likewise for the rest. At the final stage link everything together by running:

nvcc -std=c++11 -Icpp11-range -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env --cl-version 2013 -maxrregcount=0 --machine 64 -cudart static -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -I$HOME/cuShiftOr/CuShiftOr -I$HOME/cuShiftOr/CuShiftOrBenchmark -I$HOME/cuShiftOr/tinyformat -Xcompiler "-fPIC -std=c++11 -Icpp11-range -fexceptions -ffunction-sections -fdata-sections -fpermissive -Wnon-template-friend" -ccbin=/opt/cray/pe/craype/2.5.10/bin/CC -o CuShiftOrBenchmark.a   objs/CuShiftOr/device/hybrid.cu.o  objs/CuShiftOr/util/op.cu.o  objs/CuShiftOr/device/segment.cu.o  objs/CuShiftOr/host.cu.o  objs/CuShiftOr/util/timer.cu.o  objs/CuShiftOr/util/file.cu.o  objs/CuShiftOr/util/generator.cu.o  objs/CuShiftOrBenchmark/kernel.cu.o -lm

However that leads me this undefined reference issue:

objs/CuShiftOrBenchmark/kernel.cu.o: In function `void run_benchmark<unsigned int>(std::ostream&, cushiftor::device::Handler<unsigned int>*, long, unsigned int, int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)':
tmpxft_00003ea2_00000000-4_kernel.cudafe1.cpp:(.text._Z13run_benchmarkIjEvRSoPN9cushiftor6device7HandlerIT_EEljiNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE[_Z13run_benchmarkIjEvRSoPN9cushiftor6device7HandlerIT_EEljiNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE]+0x203): undefined reference to `cushiftor::util::op::bit<unsigned int>::size'
collect2: error: ld returned 1 exit status

inside op.cuh it is defined as:

template<typename type>
    struct bit {
            static const int size = sizeof(type) * 8;
    };
template<typename word>
__host__ __device__ __forceinline__ word ones(int m) {
        return ~word(0) >> (bit<word>::size - m);
}

and then on kernel.cu we have:

#include "util/op.cuh"
using cushiftor::util::op::bit;
using cushiftor::util::op::bytepack;
using cushiftor::util::op::div_up;
using cushiftor::util::op::round_up;
using cushiftor::util::op::shuffle_up;

Further more compiling with Visual Studio 2013 (which works since the project was made on that) shows some instantiations that I am not sure if g++ is doing e.g:

2>C:/Users/A/Desktop/cuShiftOr/CuShiftOrBenchmark/kernel.cu(209): warning C4244: 'argument' : conversion from 'int64_t' to 'cushiftor::env::datasize_t', possible loss of data
2>          C:/Users/A/Desktop/cuShiftOr/CuShiftOrBenchmark/kernel.cu(230) : see reference to function template instantiation 'void benchmark_with_file<word>(cushiftor::device::Handler<word> *)' being compiled
2>          with
2>          [
2>              word=unsigned int
2>          ]
2>C:/Users/A/Desktop/cuShiftOr/CuShiftOrBenchmark/kernel.cu(86): warning C4018: '<=' : signed/unsigned mismatch
2>          C:/Users/A/Desktop/cuShiftOr/CuShiftOrBenchmark/kernel.cu(132) : see reference to function template instantiation 'void run_benchmark<word>(std::ostream &,cushiftor::device::Handler<word> *,int64_t,cushiftor::env::datasize_t,int,std::string)' being compiled
2>          with
2>          [
2>              word=unsigned int
2>          ]

Now my C++11 template knowledge is very limited but I did try to instantiate also inside the kernel.cu by adding:

template class cushiftor::util::op::bit<unsigned int>;

However to no avail... Any tips are very welcomed!

Answer 1

I believe this question is essentially a duplicate of this one. However simply marking it as such with no explanation may be a bit obscure. Therefore I'm going to provide a CW answer in the hopes that someone may improve it or correct me if I am wrong.

A possible MCVE derived from the code is as follows (note that I'm converting to using g++ directly rather than nvcc, as the manifestation and workarounds are the same):

$ cat test.cpp
#include <iostream>

namespace cushiftor {
namespace util {
namespace op {

template<typename type>
struct bit {
        static const int size = sizeof(type) * 8;
};

}
}
}

#ifdef FIX2
void f(const int data){
#else
void f(const int &data){
#endif
  std::cout << "size of type in bits is: " <<  data << std::endl;
}

using cushiftor::util::op::bit;

#ifdef FIX1
template <typename T> const int bit<T>::size;
#endif

template<typename word>
void run_benchmark() {
  f(bit<word>::size);
}

int main(){

  run_benchmark<unsigned>();
}
$ g++ -o test test.cpp
/tmp/ccCW51e3.o: In function `void run_benchmark<unsigned int>()':
test.cpp:(.text._Z13run_benchmarkIjEvv[_Z13run_benchmarkIjEvv]+0x5): undefined reference to `cushiftor::util::op::bit<unsigned int>::size'
collect2: error: ld returned 1 exit status
$ g++ -DFIX1 -o test test.cpp
$ g++ -DFIX2 -o test test.cpp
$

This MCVE was created based on inspection of relevant code in kernel.cu, tinyformat.h, and op.cuh, in the original cuShiftOr project (and paper).

We see that if we compile the code as-is, it generates a similar error to that reported in the question.

Compiling with -DFIX2 demonstrates that taking the item in question (bit::size) as a reference argument is essential to witnessing of the issue.

Compiling with -DFIX1 eliminates the issue, and I believe is the correct approach in this case.

The explanation for this issue seems to be that the compiler may treat the static const member as a compile-time constant, unless it is "used" in the code. ("Used" has a special language-specific meaning here.) If it is "used" in the code, then the compiler must be able to take the address of it (sensible in the case of a reference parameter) and to take the address of it, the class/struct definition is not sufficient. (It is apparently only a declaration even though it appears in the class/struct definition. The declaration is sufficient for usage equivalently to a compile-time constant, but is not sufficient for the case where we wish to take the address of it.)

From here:

If a static data member is of const integral or const enumeration type, its declaration in the class definition can specify a constant-initializer which shall be an integral constant expression (5.19) In that case, the member can appear in integral constant expressions. The member shall still be defined in a namespace scope if it is used in the program

I believe the FIX1 code modification satisfies the bolded requirement above.

With the above preamble, I believe it is valid to mark this as a duplicate.

With respect to the behavior of the original project on windows, I would conjecture that compilers may have leeway in enforcement of this, and may in fact provide referenceable member symbols even when the program does not explicitly define them.