in cuPrintf.cuh I find the following
// DEVICE SIDE
// External function definitions for device-side code
// Abuse of templates to simulate varargs
__device__ int cuPrintf(const char *fmt);
However when I try to declare a
device fuction in a header and put the implementation into some .cu file which I do not include I get the rror message
Error: External calls are not supported (found non-inlined call to _Z17asinf_tablelookupf)
Inlining the code works fine.
Now what? This restriction is mentioned nowhere in combination with
What am I missing?
CUDA does not have a linker on the device side, so you cannot call device functions in a different .cu file (as you’ve just found out).
@tera: 1) Where in the documentation does it say that there is no linker on the device side?
2) I use cuPrintf, but I only need to include cuPrintf.cu in the routine that calls the kernel!
It’s quite hidden in Appendix E.1 of the Programming Guide:
The same statement for variables is slightly easier to find in Appendix B.2.5 “Restrictions”:
device[/font], [font=“Courier New”] shared[/font] and [font=“Courier New”] constant[/font] variables cannot be defined as external using the [font=“Courier New”]extern[/font] keyword. The only exception is for dynamically allocated [font=“Courier New”] shared[/font] variables as described in Section B.2.3.
Ok, I figured out myself that external calls do not work.
Still I would like to quote from the cuPrintf Readme (in addition to the code snippet comments you find in my original post):
The cuPrintf package consists of two device functions (i.e. called from within a CUDA kernel) and three host functions (i.e. called from within the host application). These are packaged in a single cuPrintf.cu file, along with declarations included in a separate cuPrintf.cuh header file. To use cuPrintf in your application, you must do one of the following:
a) Either: Include the header-file cuPrintf.cuh at the top of your device code, and add cuPrintf.cu to your makefile or build command-line so that the file is included in your program.
B) Or: Directly â€œ#include cuPrintf.cuâ€ at the top of your device code. In this case you should not add this file to your makefile/build-command, and you should take care to only include it once in your entire project.
Following the answers I got, option (a) is impossible ?!
It’s hard to believe, why should it be in the readme?
Or is it a matter of operating system (I have LINUX)?