What exactly is "__eh_curr_region" ?


I know this error has been discussed before, but my question concerns the actual causes of this error. I am currently porting a large chunk of Fortran code to CUDA. As soon as the last syntax error was removed I thought everything would be fine, but no! I got 5 new errors, all of them

fesft.cu(345): error: identifier “__eh_curr_region” is undefined

I ran a search on it in the forum (and on google) and it turned out that this error usually occurs when someone tries to allocate a chunk of shared memory that is not a constant size. I do not do that. Actually I do not fiddle with the memory hierarchies at all so far. This error occurs on the closing brackets of for-loops. I could not find any reference of this error or what “__eh_curr_region” is at all in the documentation. Can anybody help me on this?


I am having the same problem - porting a Fortran code, and receiving this error (although mine is being reported on the last line of an ‘if’ statement), even though I am not using anything other than device memory at the moment.

Based on the other posts where it was caused by having shared arrays of variable size, I looked at the arrays in my code. There is a local 2D array (defined within the kernel) of variable size, so I tried changing that to be of fixed size to see if that was the problem. It seemed to solve it, so it seems variable-sized arrays are not allowed in device memory either.

I just had this yesterday. The code compiled fine in Windows/VS but failed in linux.

It is indeed some shared memory array, which nvcc seems to not resolve its size correctly.

I think the fastest way to understand what’s causing this is to do this:

  • search in the offeneding kernel for the shared arrays

  • if you have an array defined like this:

    shared int smArr[SOME_SIZE]

    just change it to smArr[ 100 ] and compile

  • do this till you find the offending line.

In my code I had the following:

<in a different .cuh file>

#define VAL 20

#define VAL1 __max( VAL, 7 )


__global__ void myKernel(...)


  __shared__ int smArr[ VAL ];   // That was ok

  __shared__ int smArr1[ VAL1 ];   // That failed compilation..


hope that helps…