Problem Including in Multiple files

[GoogleCodeExporter]

Very nice performance with your sorting function!  A 2x boost on all sorting is 
always welcome.  Just looked at the updates and see that you guys have been 
very busy, so I guess this is already on the radar.  I wanted to use it in 
multiple files that are compiled separately, but received linking errors.   

What steps will reproduce the problem?
1. Include the radixsort_api.cu in different files and try to combine them at 
linking time.
   nvcc -c file1.cu
   nvcc -c file2.cu
   nvcc -c main.cu
   nvcc -o main main.o file1.o file2.o

This gives linking errors, when some variables are defined multiple times.
I attached the files, compiler_error.txt with the error I received, and the 
source files. 
Thanks,
Scott

Original issue reported on code.google.com by scott.ro...@gmail.com on 11 Aug 2010 at 11:58

Attachments:

• compile_error.txt
• b40c-link-error.tgz

[GoogleCodeExporter]

I modified a few things to get it to work.
In b40c_kernel_utils.cu, line 278:
template <int Value>
__global__ void FlushKernel()
{
}

This function is called twice in radixsort_api.cu, lines 470, 502:
FlushKernel<0> <<<_grid_size, RADIXSORT_THREADS, 
scan_scatter_attrs.sharedSizeBytes>>>();
FlushKernel<0><<<_grid_size, RADIXSORT_THREADS, 
scan_scatter_attrs.sharedSizeBytes>>>();

In file kernel/radixsort_kernel_common.cu, I changed the IsKeysOnly function to 
specialize using a class template, just after line 130.

In file kernel/radixsort_spine_kernel.cu, line 116.  I made it a template 
function.  This isn't the best way since it could be given incorrect template 
arguments, but it does compile.

In file radixsort_api.cu, to accomodate the templateness of the spine kernel, 
line 313 had to be changed.
cudaFuncGetAttributes(&_spine_scan_kernel_attrs, SrtsScanSpine<unsigned int>);

With these changes I've successfully able to compile multiple objects and then 
link them using the example files above.

Original comment by scott.ro...@gmail.com on 12 Aug 2010 at 3:24

Attachments:

• b40c_kernel_utils.cu
• radixsort_api.cu
• radixsort_kernel_common.cu
• radixsort_spine_kernel.cu

[GoogleCodeExporter]

Hi Scott,

Thanks for the report.  I encountered that problem yesterday as well when 
integrating Duane's code into Thrust.  Can you confirm that r201 resolves the 
problem?

Original comment by wnbell on 12 Aug 2010 at 11:16

• Changed state: Accepted

[GoogleCodeExporter]

Hey Nathan, 
just one more multiple definition:
line 99 in radixsort_api.cu,
bool RADIXSORT_DEBUG = false;

making it const is one fix. 

Original comment by scott.ro...@gmail.com on 13 Aug 2010 at 4:42

[GoogleCodeExporter]

Ah, I forgot about that one :)

Should be resolved in r206

Original comment by wnbell on 13 Aug 2010 at 9:30

• Changed state: Fixed