0

If I write the following CUDA code:

#include <stdio.h>

template <unsigned N>
__global__ void foo()
{
    printf("In kernel foo() with N = %u\n", N);
    if (N < 10) { return; }
    printf("Wow, N is really high!\n");
    /* a whole lot of code here which I don't want to indent */
}

int main() {
    foo<5><<<1,1>>>();
    foo<20><<<1,1>>>();
    return 0;
}

I get a compiler warning:

a.cu(8): warning: statement is unreachable
          detected during instantiation of "void foo<N>() [with N=5U]" 
(12): here

I "feel" I shouldn't be getting this warning, since the unreachable code is only unreachable for certain values of the template parameter. And if I write the "CPU equivalent", so to speak:

#include <cstdio>

template <unsigned N>
void foo()
{
    std::printf("In kernel foo() with N = %u\n", N);
    if (N < 10) { return; }
    std::printf("Wow, N is really high!\n");
    /* a whole lot of code here which I don't want to indent */
}

int main() {
    foo<5>();
    foo<20>();
    return 0;
}

and build this with gcc (5.4.0) - I don't get any warnings, even if I compile with -Wall.

Now, I can circumvent this by writing

if (not (N < 10)) { 
    printf("Wow, N is really high!\n");
    /* a whole lot of code here which I don't want to indent */
}

but I would rather avoid having to reverse my logic to jump through nvcc's "hoop". I could also write

if (not (N < 10)) { 
    return;
}
else {
    printf("Wow, N is really high!\n");
    /* a whole lot of code here which I don't want to indent */
}

but - I don't want to indent all that code (and the same problem may occur again, requiring even more indentation inside the else block.

Is there something I could do? Also, isn't this a "bug", or a misfeature I should report as a bug?

einpoklum
  • 118,144
  • 57
  • 340
  • 684
  • 2
    first thing i'd consider is adding an `else` in there – Captain Obvlious Apr 01 '17 at 15:41
  • @CaptainObvlious: See edit. – einpoklum Apr 01 '17 at 17:37
  • 1
    I don't understand what you are complaining about. You are getting a warning for the instances where the code is unreachable. Why shouldn't the compiler do that? – talonmies Apr 01 '17 at 19:48
  • 1. Because the code _is_ reachable - for certain values of the template parameters. It not being reachable in some situations is not a problem which merits a warning 2. Because other compilers (gcc, clang) realize this and do not emit a warning. – einpoklum Apr 01 '17 at 20:00
  • 2
    I am with @talonmies. If you call `foo<5>()` you are calling a function which has been "compiled" (not sure whether that is the right term for templated functions) for the special case of N==5. This result, which is different from the result for N==20, assumes N==5, without treating it as a variable. This I consider the idea of template programming. A compiler is quite right to mention code which is unreachable with that constant. Other compilers are also OK not to warn; they could implement very similar code, but be aware of the changing nature of the constant N. – Yunnosch Apr 01 '17 at 20:07
  • @Yunnosch: You're trying to rationalize an nvcc bug. Do you really believe nVIDIA decided that the GCC and clang teams are wrong about this and they will force these warnings onto people using nvcc since nVIDIA knows best? – einpoklum Apr 01 '17 at 20:09
  • 3
    Seems like ranting now. Why not file a bug at developer.nvidia.com? There's any number of plausible explanations for this behavior that don't depend on NVIDIA being a nefarious organization. It's possible that the bifurcated nature of host/device compilation paths in `nvcc` makes it more difficult to handle this case the way you would like, as compared to how host compilers handle it. I'm not a language expert so I have no idea if this is compliant or not. The warning itself is emanating from the `cudafe` tool, and is affected by that tool's `-t` switch which affects template instantiation. – Robert Crovella Apr 01 '17 at 20:17
  • @einpoklum Isn't that your question? Whether it is a bug, which should be reported, or not? I tried to answer that question. I did so based on my opinion, having the impression that you asked for it. I am also on your side, the warning is annoying. In my experience however, compiler suppliers are not very fast with fixing "errors" if they find a reasoning that it is not objectively an error from all points of view. Generally, not thinking of a certain one... – Yunnosch Apr 01 '17 at 20:31
  • @RobertCrovella: I'll do that. I had thought maybe there's some trick, or `#pragma`, or something I've missed which would let me avoid this warning easily. – einpoklum Apr 01 '17 at 20:40
  • @RobertCrovella: Correction: Will do that when the bug system stops telling me "access denied". – einpoklum Apr 08 '17 at 08:19
  • 1
    The access denied message arises out of an aggressive spam/threat detection system. My suggestion is to provide a super simple trivial report first - strip your report down to almost zero text entry of any kind. Then add to the report as necessary, after it is logged by the system. Apologies. – Robert Crovella Apr 09 '17 at 08:19
  • @RobertCrovella: First part of your advice works, second part apparently hits the same spam filter. Eventually I gave up on posting any code on the bug page and just linked back to here. That passed the filter. Thanks, though. – einpoklum Apr 09 '17 at 08:46
  • Again, apologies. If you have content you would like me to add to the bug, and give me the bug number reported by the system, I will do so. For example, a link to this question on SO would probably give nearly all the background needed. If you managed to get the link in, the description is probably sufficient. If not, and you want me to add it, just let me know the bug number. – Robert Crovella Apr 09 '17 at 08:56
  • @RobertCrovella: I linked here, so it's fine. – einpoklum Apr 09 '17 at 09:34

1 Answers1

1

What about:

template<unsigned N, bool>
struct FooImpl
{
  static void foo()
  {
    std::printf("In kernel foo() with N = %u\n", N);
  }
};

template<unsigned N>
struct FooImpl<N, false>
{
  static void foo()
  {
    std::printf("In kernel foo() with N = %u\n", N);
    std::printf("Wow, N is really high!\n");
    /* a whole lot of code here which I don't want to indent */
  }
};

template <unsigned N>
__global__ void foo()
{
  FooImpl<N, N < 10>::foo();
}
Tomek
  • 4,554
  • 1
  • 19
  • 19
  • Well, that is "another way" to achieve the same, but obviously we should prefer using real code rather than template metaprogramming to effect conditions... this way is quite a lot more code, and more confusing to readers. Also, again, if I have nested `if (condition) { return; }` lines in my original code, I would need to nest +Impl functions - not fun... so it's not quite what I was hoping for. – einpoklum Apr 01 '17 at 20:02