nvcc: Compilation error in case of protected ctor

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi Everyone,

I ran into a weird compilation error which I cannot get over. I have a class with a member function defined in a separate .cu file. This .cu file includes some headers. In one of these headers there is a class defined which publicly inherits from a base class with a protected ctor (so that the base class cannot be instantiated). nvcc complains about the protected function, throwing an error.
Although, excluding the .cu file from the compilation by defining the member function in a .cpp file which still includes the same headers, compiling with g++ which was previously set as the host compiler for nvcc, compilation succeeds.

(Sorry about the long post but the example code is actually really simple.)

Here is a dummy example resulting in the error:

MainClass.h

class MainClass{
	public:
		void kernelCallerMember();
};

cudaFile.cu

#include "MainClass.h"

#include "Bar.h"

__global__ void dummyKernel(){

}

void MainClass::kernelCallerMember(){
	dim3 blocksPerGrid{256, 256};
	dim3 threadsPerBlock{32, 32};
	dummyKernel<<<blocksPerGrid, threadsPerBlock>>>();
}

Bar.h

#include "Foo.h"

class Bar : public Foo{
	public:
		Bar(int _a) : Foo{_a}{}
};

Foo.h

class Foo{
	protected:
		Foo(int _a) : a{_a}{}

		int a;
};

The command I run is:

nvcc cudaFile.cu -c -std=c++11 -ccbin=g++

The resulting error:

nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
Bar.h(5): error: protected function "Foo::Foo(int)"
Foo.h(3): here is not accessible through a "Foo" pointer or object

1 error detected in the compilation of "/tmp/tmpxft_00004c46_00000000-7_cudaFile.cpp1.ii".

Proof that this would be no problem for g++:

cppFile.cpp

#include "MainClass.h"

#include "Bar.h"

void MainClass::kernelCallerMember(){

}

Command to compile:

g++ cppFile.cpp -c -std=c++11

Result: successful compilation

Some notes:

  • Running nvcc with the --dryrun flag, then running the commands one by one, it looks that the first gcc preprocess succeeds, but then cudafe gives the above error
  • The same code compiles in a cross-compiling environment, where the host compiler compiles to the TX1's 64-bit ARM architecture (while the error happens when compiling for x86_64 ubuntu)

Any help is much appreciated!

Thanks,
Daniel

2

It’s evident that you are using CUDA 8. Try CUDA 9. This compiles cleanly for me in CUDA 9:

$ cat t14.cu
class MainClass{
        public:
                void kernelCallerMember();
};


class Foo{
        protected:
                Foo(int _a) : a{_a}{}

                int a;
};

class Bar : public Foo{
        public:
                Bar(int _a) : Foo{_a}{}
};

__global__ void dummyKernel(){

}

void MainClass::kernelCallerMember(){
        dim3 blocksPerGrid{256, 256};
        dim3 threadsPerBlock{32, 32};
        dummyKernel<<<blocksPerGrid, threadsPerBlock>>>();
}
$ nvcc -c t14.cu -std=c++11
$
3

Hi txbox,

Thank you for your reply! I tried it using CUDA 9, and it compiles fine indeed.