CUDA 8.0 Visual Studio 2012 linker options when using Relocatable Device Code

Question

I have a class-based C++ code similar to this and have followed the structure for declaring a __device__ method of the class. I have also gone to Project Properties-> CUDA C/C++ -> Common -> Generate Relocatable Device Code and changed it to YES but here it says the linking options need modifying and I do not understand how to do this in the VS options from this information. How do I do this to get it to link?

**Removed VC studio settings in original post as not relevant**

Test code is:

kernel.cu

#include "Test.cuh"

// CUDA kernel
__global__ void myKernel(Test* t) {

// Call device function
t->device_function();
}


// Entry point
int main()
{

// Create object (host and device copies)
Test* t = new Test();
Test* t_dev = NULL;
cudaMalloc(&t_dev, sizeof(Test));

// Copy object
cudaMemcpy(t_dev,t,sizeof(Test),cudaMemcpyHostToDevice);

// Call kernel passing pointer to device copy of the object
myKernel<<<1,1>>>(t_dev);

}

Test.cuh

class Test
{
public:
Test(void);
~Test(void);

__device__ void device_function();
};

Test.cu

#include "Test.cuh"
Test::Test(void){}
Test::~Test(void){}

__device__ void Test::device_function() {

// Do something

}

EDIT

Adding compilation log as I think it might not be compiling the device implementation properly.

1> C:\C++ Projects\CudaTest\CudaTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\bin\nvcc.exe" -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin"  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include"     --keep-dir Release -maxrregcount=0  --machine 32 --compile      -DWIN32 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /Zi  /MD " -o Release\kernel.cu.obj "C:\C++ Projects\CudaTest\CudaTest\kernel.cu" -clean
1>  kernel.cu
1>  Compiling CUDA source file kernel.cu...
1>  
1>  C:\C++ Projects\CudaTest\CudaTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\bin\nvcc.exe" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin" -rdc=true -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include"     --keep-dir Release -maxrregcount=0  --machine 32 --compile -cudart static     -DWIN32 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /Zi  /MD " -o Release\kernel.cu.obj "C:\C++ Projects\CudaTest\CudaTest\kernel.cu" 
1>  kernel.cu
1>  Test.cu
1>  
1>  C:\C++ Projects\CudaTest\CudaTest>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\bin\nvcc.exe" -dlink -o Release\CudaTest.device-link.obj -Xcompiler "/EHsc /W3 /nologo /O2 /Zi  /MD " -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\Win32" cudart.lib kernel32.lib user32.lib gdi32.lib winspool.lib comdlg32.lib advapi32.lib shell32.lib ole32.lib oleaut32.lib uuid.lib odbc32.lib odbccp32.lib  -gencode=arch=compute_20,code=sm_20  --machine 32 Release\kernel.cu.obj 
1> CUDALINK : nvlink error : Undefined reference to '_ZN4Test15device_functionEv' in 'Release/kernel.cu.obj'`

`Test.cpp` can't possibly be compiling. Could you paste the actual compilation log? The cut and paste from VC you have posted is completely unhelpful in understand what might be going wrong. — talonmies, Jul 13 '16 at 11:10
Is this something to do with only `kernel.cu` being compiled by the CUDA compiler and not the `Test` class so the device code is not embedded correctly? -- Sorry , just wrote this comment without having seen yours. I guess I'm thinking along the right lines then. — CodingLumis, Jul 13 '16 at 11:23
@talonmies OK, I think I have the problem -- sorry I've been editing since I've posted which has not helped you. I renamed my headers and source files to the cuda extensions thinking that would force the `nvcc` compiler to compile the relevant code in those files (i.e. the device function) but it does not and the error remained. If I added the files fresh in VS then it works. Do you know if it is possible to convert the files without having to add them to the project again? — CodingLumis, Jul 13 '16 at 11:39
I seem to be good at answering my own questions while waiting for an answer here... If you rename the files, you need to edit the `*.vcxproj` file and under `<\ItemDefinitionGroup>` add — CodingLumis, Jul 13 '16 at 11:44
I don't use VS, so I don't know how to do that, I'm afraid. The original issue was caused by `Test.{cu,cpp}` not being included in the device linking stage of the build. — talonmies, Jul 13 '16 at 11:46
There are various CUDA sample projects that include all the necessary settings for projects that use relocatable device code with device linking. You could also start by studying one of those. You should never have to be editing a vcxproj file directly. If you want to add a CUDA source file to your project, use the VC method to do this (right click in project explorer, etc.) and choose to add it, then either select the file you already have, or paste the contents into your newly added CUDA source file. Then just remove the C++ source file from your project. — Robert Crovella, Jul 13 '16 at 13:53

CodingLumis · Accepted Answer · 2016-12-01T17:44:04.787

OK, so VS with the CUDA plug-in has two different Item Types that it catalogues differently, a CUDA C/C++ source/header file and a C/C++ source/header and you cannot simply convert one to another by renaming *.cpp to *.cu which is what I was doing originally to try force it to compile the class into an object which I could link to.

I re-added the files Test.cpp and Test.h with exactly the same source code as Test.cu and Test.cuh respectively by right-clicking in the solution explorer and going to Add Item -> CUDA source/header file and now it builds the kernel.cu and the Test.cu/h into two objects which are both supplied during the linking stage and the problem disappears.

Thanks to @talonmies for helping point me in the right direction.

EDIT

As I mentioned in the comments, you can simply rename the files if you still need to tell VS that they are now a different Item Type. Originally I managed to do this by editing the *.vcxproj file to tell the CUDA compiler to compile the renamed files. This is done by changing the <ItemGroup> definitions like this:

<ItemGroup>
  <CudaCompile Include="main.cu" />
  <CudaCompile Include="RenamedSrcFile.cu" />
</ItemGroup>

<ItemGroup>
  <ClInclude Include="RenamedHeaderFile.cuh" />
  <ClInclude Include="myOtherHeader.h" />
</ItemGroup>

This line <CudaCompile Include="RenamedSrcFile.cu" /> is the key and includes this file in the CUDA compile stage. Regular C++ files appear under the group <ClCompile> although none are given in this example.

As indicated by Robert's comment you should never need to edit the *.vcxproj file directly as VS will manage the file for you based on the options you select in VS. By (re-)adding the files as I did earlier, this swap is taken care of by VS.

EDIT2

As per Drop's suggestion in the comments, you may achieve this swap without re-adding files by right-clicking on the file in the solution explorer and under Properties->General you may select the Item Type to be a CUDA C/C++ rather than C/C++ to catalogue it correctly.

If deduced compiler for a source file doesn't match your expectations, you could simply right click on a source file (or a bunch of selected files), go to Properties and chose between C++ and CUDA compiler (by choosing appropriate "type"). Note, that in this case, you don't need to worry about header files, as they aren't compiled directly (and also have a separate "type" in Visual Studio) — Ivan Aksamentov - Drop, Jul 13 '16 at 23:56

CUDA 8.0 Visual Studio 2012 linker options when using Relocatable Device Code

1 Answers1