cicc generating unecessary warnings

Accelerated Computing CUDA CUDA Programming and Performance
1

I’m having an issue where cicc is generating a flood of warnings I neither want nor need, and I can’t seem to be able to disable them. I do not want to disable warnings across the board, but I really would like a way to make this particular one be silent.

In my real code I’m getting pages of a particular warning so it’s not something I can easily ignore, but I trimmed things down to a minimal working example to demonstrate my issue:

namespace {

template <int N>
struct Foo {
   __device__ __host__ void A() {}
   __device__ __host__ void B() {}
   __device__ __host__ void C() {}
};

}

__global__ void dummy()
{
    Foo<1> one;
    one.A();
    one.B();
    one.C();

    Foo<2> two;
    two.A();
}

int main()
{
    Foo<1> one;
    one.A();
    one.B();
    one.C();

    Foo<2> two;
    two.B();

    dummy<<<1,1>>>();
}

If I compile this code I get the warning: function “::Foo::C [with N=2]” was declared but never referenced.

The class Foo has three functions all marked both host and device so that I could show that using the function either on the host or on the device (but not necessarily both) is enough to avoid the warning on a function.

From what I can tell these are the important bits:

  • The warning only appears while I'm using an anonymous namespace. The anonymous namespace gives functions static linkage, which is something that will trigger the "unused-functions" warning in gcc if it was enabled directly or "-Wall" was used. However I did not enable this warning, nvcc seems to be doing that itself.
  • It's not even clear to me that this warning should apply when the functions are members of a class. At the least I cannot get g++ with similar code to trigger this warning
  • Since the class is templated and one instance of the template actually has all the functions invoked, the warning seems inappropriate. The function is used and commenting it out will cause a compiler error

    • I know that you can pass arguments to the backend compiler directly, but using “-Xcompiler -Wno-unused-function” does not fix the issue. If I run nvcc with verbose output it seems that cicc itself is issuing this warning.

    nvcc tmp.cu -Xcompiler -Wno-unused-function --verbose
#$ _SPACE_= 
#$ _CUDART_=cudart
#$ _HERE_=/usr/local/cuda/bin
#$ _THERE_=/usr/local/cuda/bin
#$ _TARGET_SIZE_=
#$ _TARGET_DIR_=
#$ _TARGET_SIZE_=64
#$ TOP=/usr/local/cuda/bin/..
#$ NVVMIR_LIBRARY_DIR=/usr/local/cuda/bin/../nvvm/libdevice
#$ LD_LIBRARY_PATH=/usr/local/cuda/bin/../lib:/opt/rh/llvm-toolset-7/root/usr/lib64:/opt/rh/devtoolset-6/root/usr/lib64:/opt/rh/devtoolset-6/root/usr/lib
#$ PATH=/usr/local/cuda/bin/../nvvm/bin:/usr/local/cuda/bin:/home/UNIXHOME/bbyington/.local/bin:/home/UNIXHOME/bbyington/bin:/home/UNIXHOME/bbyington/bin/bin:/home/UNIXHOME/bbyington/rtags/bin:/home/UNIXHOME/bbyington/emacs/bin/:/home/UNIXHOME/bbyington/rtags/bin:/home/UNIXHOME/bbyington/emacs/bin/:/home/UNIXHOME/bbyington/.local/bin:/home/UNIXHOME/bbyington/bin:/home/UNIXHOME/bbyington/bin/bin:/home/UNIXHOME/bbyington/rtags/bin:/home/UNIXHOME/bbyington/emacs/bin/:/opt/rh/llvm-toolset-7/root/usr/bin:/opt/rh/llvm-toolset-7/root/usr/sbin:/opt/rh/devtoolset-6/root/usr/bin:/usr/lib64/qt-3.3/bin:/usr/local/bin:/usr/local/sbin:/usr/bin:/usr/sbin:/bin:/sbin:/usr/local/cuda/bin:/usr/local/cuda/bin:/usr/local/cuda/bin:/usr/local/cuda/bin:/usr/local/cuda/bin
#$ INCLUDES="-I/usr/local/cuda/bin/..//include"  
#$ LIBRARIES=  "-L/usr/local/cuda/bin/..//lib64/stubs" "-L/usr/local/cuda/bin/..//lib64"
#$ CUDAFE_FLAGS=
#$ PTXAS_FLAGS=
#$ gcc -std=c++14 -D__CUDA_ARCH__=300 -E -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS -D__CUDACC__ -D__NVCC__  -Wno-unused-function "-I/usr/local/cuda/bin/..//include"    -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=0 -D__CUDACC_VER_BUILD__=130 -include "cuda_runtime.h" -m64 "tmp.cu" > "/tmp/tmpxft_00007b20_00000000-8_tmp.cpp1.ii" 
#$ cicc --c++14 --gnu_version=60301 --allow_managed   -arch compute_30 -m64 -ftz=0 -prec_div=1 -prec_sqrt=1 -fmad=1 --include_file_name "tmpxft_00007b20_00000000-2_tmp.fatbin.c" -tused -nvvmir-library "/usr/local/cuda/bin/../nvvm/libdevice/libdevice.10.bc" --gen_module_id_file --module_id_file_name "/tmp/tmpxft_00007b20_00000000-3_tmp.module_id" --orig_src_file_name "tmp.cu" --gen_c_file_name "/tmp/tmpxft_00007b20_00000000-5_tmp.cudafe1.c" --stub_file_name "/tmp/tmpxft_00007b20_00000000-5_tmp.cudafe1.stub.c" --gen_device_file_name "/tmp/tmpxft_00007b20_00000000-5_tmp.cudafe1.gpu"  "/tmp/tmpxft_00007b20_00000000-8_tmp.cpp1.ii" -o "/tmp/tmpxft_00007b20_00000000-5_tmp.ptx"
tmp.cu(7): warning: function "<unnamed>::Foo<N>::C [with N=2]" was declared but never referenced

#$ ptxas -arch=sm_30 -m64  "/tmp/tmpxft_00007b20_00000000-5_tmp.ptx"  -o "/tmp/tmpxft_00007b20_00000000-9_tmp.sm_30.cubin" 
#$ fatbinary --create="/tmp/tmpxft_00007b20_00000000-2_tmp.fatbin" -64 "--image=profile=sm_30,file=/tmp/tmpxft_00007b20_00000000-9_tmp.sm_30.cubin" "--image=profile=compute_30,file=/tmp/tmpxft_00007b20_00000000-5_tmp.ptx" --embedded-fatbin="/tmp/tmpxft_00007b20_00000000-2_tmp.fatbin.c" --cuda
#$ rm /tmp/tmpxft_00007b20_00000000-2_tmp.fatbin
#$ gcc -std=c++14 -E -x c++ -D__CUDACC__ -D__NVCC__  -Wno-unused-function "-I/usr/local/cuda/bin/..//include"    -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=0 -D__CUDACC_VER_BUILD__=130 -include "cuda_runtime.h" -m64 "tmp.cu" > "/tmp/tmpxft_00007b20_00000000-4_tmp.cpp4.ii" 
#$ cudafe++ --c++14 --gnu_version=60301 --allow_managed  --m64 --parse_templates --gen_c_file_name "/tmp/tmpxft_00007b20_00000000-5_tmp.cudafe1.cpp" --stub_file_name "tmpxft_00007b20_00000000-5_tmp.cudafe1.stub.c" --module_id_file_name "/tmp/tmpxft_00007b20_00000000-3_tmp.module_id" "/tmp/tmpxft_00007b20_00000000-4_tmp.cpp4.ii" 
#$ gcc -std=c++14 -D__CUDA_ARCH__=300 -c -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS -Wno-unused-function "-I/usr/local/cuda/bin/..//include"   -m64 -o "/tmp/tmpxft_00007b20_00000000-10_tmp.o" "/tmp/tmpxft_00007b20_00000000-5_tmp.cudafe1.cpp" 
#$ nvlink --arch=sm_30 --register-link-binaries="/tmp/tmpxft_00007b20_00000000-6_a_dlink.reg.c"  -m64   "-L/usr/local/cuda/bin/..//lib64/stubs" "-L/usr/local/cuda/bin/..//lib64" -cpu-arch=X86_64 "/tmp/tmpxft_00007b20_00000000-10_tmp.o"  -o "/tmp/tmpxft_00007b20_00000000-11_a_dlink.sm_30.cubin"
#$ fatbinary --create="/tmp/tmpxft_00007b20_00000000-7_a_dlink.fatbin" -64 -link "--image=profile=sm_30,file=/tmp/tmpxft_00007b20_00000000-11_a_dlink.sm_30.cubin" --embedded-fatbin="/tmp/tmpxft_00007b20_00000000-7_a_dlink.fatbin.c" 
#$ rm /tmp/tmpxft_00007b20_00000000-7_a_dlink.fatbin
#$ gcc -std=c++14 -c -x c++ -DFATBINFILE="\"/tmp/tmpxft_00007b20_00000000-7_a_dlink.fatbin.c\"" -DREGISTERLINKBINARYFILE="\"/tmp/tmpxft_00007b20_00000000-6_a_dlink.reg.c\"" -I. -D__NV_EXTRA_INITIALIZATION= -D__NV_EXTRA_FINALIZATION= -D__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__  -Wno-unused-function "-I/usr/local/cuda/bin/..//include"    -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=0 -D__CUDACC_VER_BUILD__=130 -m64 -o "/tmp/tmpxft_00007b20_00000000-12_a_dlink.o" "/usr/local/cuda/bin/crt/link.stub" 
#$ g++ -Wno-unused-function -m64 -o "a.out" -std=c++14 -Wl,--start-group "/tmp/tmpxft_00007b20_00000000-12_a_dlink.o" "/tmp/tmpxft_00007b20_00000000-10_tmp.o"   "-L/usr/local/cuda/bin/..//lib64/stubs" "-L/usr/local/cuda/bin/..//lib64" -lcudadevrt  -lcudart_static  -lrt -lpthread  -ldl  -Wl,--end-group
    2

    my suggestion would be to file a bug at developer.nvidia.com

    3

    I’ll be glad too. Does this mean there is no way to suppress warnings coming from cicc? I tried poking at that executable manually but it doesn’t seem to have any help message so I can see what options it may accept.

    4

    Nothing is formally documented for it AFAIK.

    If you want to experiment, this may have some things to try:

    [url]cuda - How to disable a specific nvcc compiler warnings - Stack Overflow

    I don’t have an exact answer for how to do it, so if nothing there is useful then I have no further suggestions.

    5

    OK. Thank you very much for the information.