`printf(...)` with "%s" doesn't work in larger format string, but does when not

Accelerated Computing CUDA CUDA Programming and Performance
1

I have the following CUDA code:

#if 1 //Doesn't print anything
	printf("(%s:%d): ",filename,line);
#else //Prints, but bloated code and interleaved output
	printf("(");
	printf("%s",filename);
	printf(":%d): ",line);
#endif

The one-liner prints nothing. The second one works, but output from multiple threads becomes interleaved. I conclude that something is broken with the "%s" specifier implementation.

What’s going on? Can this be fixed and/or worked around?

Additional details that may be relevant but probably aren’t:

  • Windows 10, latest patches
  • GTX 1080, driver 435 (bug in 3rd-party library prevents testing 436)
  • In context of OptiX 7.0.0 application
    • 2

    I don’t seem to have any difficulty with it:

    $ cat t1517.cu
#include <stdio.h>

__host__ __device__ void f(){
        const char *filename = "hello";
        int line = 4;
#if 1 //Doesn't print anything
        printf("(%s:%d): ",filename,line);
#else //Prints, but bloated code and interleaved output
        printf("(");
        printf("%s",filename);
        printf(":%d): ",line);
#endif
}

__global__ void k(){

  f();
}

int main(){

  printf("\nCPU:\n");
  f();
  printf("\nGPU:\n");
  k<<<1,1>>>();
  cudaDeviceSynchronize();
  printf("\n");
  return 0;
}
$ nvcc -o t1517 t1517.cu
$ cuda-memcheck ./t1517
========= CUDA-MEMCHECK

CPU:
(hello:4):
GPU:
(hello:4):
========= ERROR SUMMARY: 0 errors
    3

    I tried to reproduce the problem in a simple OptiX application that does not depend on the mentioned 3rd-party library (and so runs in 436), and I was not able to reproduce the problem there either.

    Either this issue was already fixed in 436, or it is something more-complicated related to my real application that I will need to isolate once I can run it in 436. If it turns out to be the latter, I will update here.

    Thanks,

    4

    Is this the only printf? The print fifo is circular and relatively small.(Search on cudaLimitPrintfFifoSize.)