Unable to compile with nvcc, gcc-4.3.3 support for exception handling is disabled

When compiling with nvcc, I get a lot of these:
/usr/lib/gcc/i686-pc-linux-gnu/4.3.2/include/g+±v4/bits/locale_classes.h(407): error: support for exception handling is disabled
spread out over a bunch of gcc headers.

While I understand that “exception handling is disabled”, I don’t understand why I get the message.
Is there a simple workaround?

[edit]The same code compiled fine with cuda 2.1, and doesn’t compile either emulation or “normal” mode.[/edit]

My bad, only emulation fails.
Am I to understand from
that I can’t expect it to work with 2.2 at all?

Correctamundo. You should stop relying on emulation anyway because it’s useless outside of valgrind support, really, and just use cuda-gdb instead if you’re actually debugging.

cudagdb plus ddt (plus a non-gnu host compiler) wreak havok. tmurray, your statement is oversimplifying things, sorry :) Don’t get me wrong, cudagdb is a great step forward. It’s not just quite usable yet in all cases. printf()-style debugging is still, for me, the only way to go.

I must completely agree. I ran into the same problem today and this is really frustrating, because now I’m forced to rewrite a lot of code to get device emulation working… yeah!

You should probably know–I am on a crusade to kill device emulation because it is a pain (from a development standpoint and from a user’s standpoint–look at how many threads on here are “my code works in deviceemu but not on the device itself”). What you’ve said to me is “if you have printf that I can call from the GPU, that’s good enough for me.” That’s a good data point…

I saw a similar thing demonstrated on the Cell: printf on the SPU was a stub that passed the real work up to the PPU. It looked really convenient for debugging.

If you guys can come up with valgrind for the GPU and printf that runs on the device, I think that would kill off most of the use cases for device emulation. (Except for people who want to do CUDA development on a computer without a real CUDA device. In that case, the long hoped-for nvcc --multicore would be better anyway…)

I think valgrind and printf for the device is not really enough. What’s with a simple bank-conflict-checker, without device emulation you won’t get that working. And as seibert mentioned, first steps without a cuda enabled device can only be done with emulation. Of course I can understand the problems from a development standpoint, but I would sad if the emulation gets removed…

I second seibert’s statement. Personally, I’ve only used deviceemu together with valgrind, no other usecase for me (and I frequently forget that I need to link against the emu versions of cublas…). CUDA on non-CUDA hosts is not something I am doing, I’m maintaining a true CPU fallback anyway. In the long run, MCUDA is what I’ll be using (or OpenCL).

Currently it’s a bit annoying to comment out bits and pieces of my kernels, let them dump some values into global arrays, and read that back on the host. So yes, printf from kernels would be a real time-saver if it’s reasonably easy to do.

I think what tmurray is getting at (and we’ve sort of hijacked this thread now for a more general discussion) is that the current use cases for emulation could better served by other tools, and he needs to know what those use cases are so they can all be met before finally retiring device emulation. You’ve added another one to the list:

  • Debugging: cuda-gdb and a hypothetical device-side printf

  • Memory checking, bank conflict checking: a hypothetical GPU valgrind

  • Development on non-CUDA systems: an enhanced nvcc that generates efficient multithreaded & SSE enabled code so that your CUDA kernel actually runs well on a non-CUDA system. (This has already been demonstrated to be possible in a paper, and rumor is that NVIDIA may also be working on it internally.) If you are going to have to run your CUDA code on your CPU, wouldn’t it be nice if it ran decently well, unlike current emulation mode?

This list contains two hypothetical things, and one which may be in development but not finished yet. I don’t think emulation mode will be going away anytime soon, but if all of these things were achieved, then emulation mode as currently implemented would serve no purpose.

yeah, and for what it’s worth, seibert’s been prioritising things just in the same way as I’d have.

printf([%d] %f\n", blockDim.x*blockIdx.x+threadIdx.x, somevalue);

would be just great.

full valgrind on the device is probably quite challenging, it needs to provide more than the current valgrind (cachegrind? simdgrind?)

mcuda instead of deviceemu is, for me, opencl. Of course, NVIDIA implements CL on top of CUDA (as an alternative to C for CUDA aka the runtime API, probably by replacing/functionally-equivalent to the driver API). So that’s, in my opinion, decoupled from the rest of the feature requests.

I’m gonna try using cuda-gdb, but my reason for using emulation is printf. Implementing that GPU would probably completely remove my need for emulation.

I’m jumping in on this conversation from the Windows forums where I posted questions about being unable to compile for Emu after 2.2 was released. I can understand wanting to simplify the architecture and not wanting to hear complaints about Emu build works but device does not, but Windows does not currently have gdb available. I’ve had to port a lot of my code to Linux, configure my machine for dual boot, and start the painful process of relearning gdb - most of you probably hate Windows, but I hope you can admit that the Visual Studio debugger is quite good.

Anyhow, printf would satisfy a lot of my needs, but when tracking down a tough problem the emulator worked great. In my opinion the more tools you have in your toolbelt the better. Also, this move away from the emulator may alianate a lot of Windows developers. I don’t want to start a platform war, I just want to point out that the more people developing for CUDA the better for all of us already working with it (no metter the platform), and limiting the debugger may reduce adoption of this powerful architecture.


Certainly real debugging on Windows is a requirement before getting rid of emulation. Forcing people to switch operating systems to debug code is a non-solution.

Don’t worry, I assumed as much. :) Just getting more and more useful data on what’s actually necessary to get rid of it…