Is my CUDA kernel really runs on device or is being mistekenly executed by host in emulation?

I just got my GPU-enabled video card and started playing with CUDA. Just to get my head straight with blocks and threads I wrote a simple kernel that just stores its identifiers to the shared memory that I later copy back to host and print. But then I though, why not simply use printf inside the kernel function? I have tried that even though I believed that it was impossible. Here is what my attempt looked like:

__global__ void printThreadXInfo (int *data) { int i = threadIdx.x; data[i] = i; printf ("%d\n", i); }

.. but all of the sudden I saw the output in console. Then I searched developer's manual and found printf mentioned in the section about device emulation. It was said that device emulation provides a benefit of running a host-specific code in the kernel, like calling printf.

I don't really need to call printf. But now I am a little bit confused. I have two assumption. First is that NVidia developers implemented some specific printf on device that somehow transparently for the developer accesses calling process and executed standard printf function, and takes care of memory copying etc. That sounds a bit crazy. Another assumption is that the code I have compiled somehow runs in emulation rather than on a real device. But that doesn't sound right either because I simply measured a performance of adding two numbers on 1 million elements array and CUDA kernel manages to do it like 200 faster than I can do on a CPU. Or maybe it runs in emulation when it detects some host-specific code? If that is true, why am I not issued a warning then?

Please help me sort it out. I am using NVidia GeForce GTX 560 Ti on Linux (Intel Xeon, 1 CPU with 4 physical cores, 8 GB of RAM, if that matters). Here is my nvcc version:

$ /usr/local/cuda/bin/nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2011 NVIDIA Corporation
Built on Thu_May_12_11:09:45_PDT_2011
Cuda compilation tools, release 4.0, V0.2.1221

And here is how I compile my code:

/usr/local/cuda/bin/nvcc -gencode=arch=compute_20,code=\"sm_21,compute_20\" -m64 --compiler-options -fno-strict-aliasing -isystem /opt/boost_1_46_1/include -isystem /usr/local/cuda/include -I../include --compiler-bindir "/usr/local/cuda/bin" -O3 -DNDEBUG -o build_linux_release/ThreadIdxTest.cu.o -c ThreadIdxTest.cu

/usr/local/cuda/bin/nvcc -gencode=arch=compute_20,code=\"sm_21,compute_20\" -m64 --compiler-options -fno-strict-aliasing -isystem /opt/boost_1_46_1/include -isystem /usr/local/cuda/include -I../include --compiler-bindir "/usr/local/cuda/bin" -O3 -DNDEBUG --generate-dependencies ThreadIdxTest.cu | sed -e "s;ThreadIdxTest.o;build_linux_release/ThreadIdxTest.cu.o;g" > build_linux_release/ThreadIdxTest.d

g++ -pipe -m64 -ftemplate-depth-1024 -fno-strict-aliasing -fPIC -pthread -DNDEBUG -fomit-frame-pointer -momit-leaf-frame-pointer -fno-tree-pre -falign-loops -Wuninitialized -Wstrict-aliasing -ftree-vectorize -ftree-loop-linear -funroll-loops -fsched-interblock -march=native -mtune=native -g0 -O3 -ffor-scope -fuse-cxa-atexit -fvisibility-inlines-hidden -Wall -Wextra -Wreorder -Wcast-align -Winit-self -Wmissing-braces -Wmissing-include-dirs -Wswitch-enum -Wunused-parameter -Wredundant-decls -Wreturn-type -isystem /opt/boost_1_46_1/include -isystem /usr/local/cuda/include -I../include -L/opt/boost_1_46_1/lib -L/usr/local/cuda/lib64 -lcudart -lgtest -lgtest_main build_linux_release/ThreadIdxTest.cu.o ../src/build_linux_release/libspartan.a -o build_linux_release/ThreadIdxTest

... and by the way, both host code and kernel code is mixed in one source file with .cu extension (maybe I am not supposed to do that, but I saw this style in SDK examples).

Your help is highly appreciated. Thank you!


As of CUDA ?3.1?, they no longer do any device emulation. Printf's are now supported in the kernel.


  • Performance difference using entity framework in 32 bit and 64 bit
  • Slow Apache httpclient 4.1 compared to JMeter
  • RecyclerView programmatically click
  • ASP.NET Razor/Webmatrix sites logout user too quickly. How to change?
  • Most practical way to make multiple AJAX requests
  • Optimizing Memory in Android when managing Fragment Transactions
  • HttpClient with WebView
  • GPS coordinates on mobile phones
  • twisted.internet.error.ConnectError when run scrapy spider
  • Why isn't Kubernetes service DNS working?
  • TensorFlow C++, runtime issue
  • three.js WebVR example code works on threejs.org but not on my local server
  • Deployments not visible in Kubernetes Dashboard
  • Extracting individual digits from a float
  • Mongodb update() vs. findAndModify() performace
  • Disabling sound of embedded flash object with html
  • Which browser have this strange user agent? (IOS device)
  • It is possible use the same sql azure instance from two different cloud service of two different sub
  • Why isn't my “Fizz Buzz” test in R working?
  • netsh acl setting (need alternative method - registry settings?)
  • SonarQube: Cannot deactivate rule with missing quality profile
  • how does System.Web.HttpRequest::PathInfo work?
  • How do I include a SWC in an AS2 Flash project?
  • Looking for good analogy/examples for monitor verses semaphore
  • How to add a focus style to an editable ComboBox in WPF
  • Breaking out column by groups in Pandas
  • How solve “Qt: Untested Windows version 10.0 detected!”
  • ViewController With Transparent Background Entering Current ViewController With Push Transition
  • How can I speed up CURL tasks?
  • How do I superscript characters in a UIButton?
  • Time complexity of a program which involves multiple variables
  • Row Count Is Returning the incorrect number using RaptureXML
  • DomPDF {PAGE_NUM} not on first page
  • Splitting given String into two variables - php
  • NetLogo BehaviorSpace - Measure runs using reporters
  • C# - Is there a limit to the size of an httpWebRequest stream?
  • TFS: Get latest causes slow project reloading
  • Perl system calls when running as another user using sudo
  • ActionScript 2 vs ActionScript 3 performance
  • Build own AppleScript numerical error handling