CUDA Errors while running "CUDA By Example" julia_gpu.cu

Question

I am running the example julia_gpu.cu from the book CUDA By Example, using CUDA 6.0 in Visual Studio Express 2012. Here's the source code for reference:

#include <book.h>
#include <cpu_bitmap.h>

#define DIM 1000

struct cuComplex {
    float   r;
    float   i;
    cuComplex( float a, float b ) : r(a), i(b)  {}
    __device__ float magnitude2( void ) {
        return r * r + i * i;
    }
    __device__ cuComplex operator*(const cuComplex& a) {
        return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
    }
    __device__ cuComplex operator+(const cuComplex& a) {
        return cuComplex(r+a.r, i+a.i);
    }
};

__device__ int julia( int x, int y ) {
    const float scale = 1.5;
    float jx = scale * (float)(DIM/2 - x)/(DIM/2);
    float jy = scale * (float)(DIM/2 - y)/(DIM/2);

    cuComplex c(-0.8, 0.156);
    cuComplex a(jx, jy);

    int i = 0;
    for (i=0; i<200; i++) {
        a = a * a + c;
        if (a.magnitude2() > 1000)
            return 0;
    }

    return 1;
}

__global__ void kernel( unsigned char *ptr ) {
    // map from blockIdx to pixel position
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;

    // now calculate the value at that position
    int juliaValue = julia( x, y );
    ptr[offset*4 + 0] = 255 * juliaValue;
    ptr[offset*4 + 1] = 0;
    ptr[offset*4 + 2] = 0;
    ptr[offset*4 + 3] = 255;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
};

int main( void ) {
    DataBlock   data;
    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char    *dev_bitmap;

    HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap, bitmap.image_size() ) );
    data.dev_bitmap = dev_bitmap;

    dim3    grid(DIM,DIM);
    kernel<<<grid,1>>>( dev_bitmap );

    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );

    HANDLE_ERROR( cudaFree( dev_bitmap ) );

    bitmap.display_and_exit();
}

Building the above yielded a series of errors, which were all of the same form:

Error   2   error : calling a __host__ function("cuComplex::cuComplex") from a __device__ function("julia") is not allowed
Error   4   error : calling a __host__ function("cuComplex::cuComplex") from a __device__ function("cuComplex::operator *") is not allowed
Error   5   error : calling a __host__ function("cuComplex::cuComplex") from a __device__ function("cuComplex::operator +") is not allowed

...and so on.

I am having trouble understanding what the problem is, since as far as I was aware, __device__ functions (like julia) are free to call and create struct objects (like cuComplex). What exactly is the problem here? Is the code provided in CUDA By Example (which was written back in 2010) broken by some of the changes introduced in the more recent CUDA 6.0 update?

Answer 1

I have already submitted errata to CUDA team some time ago (without any response yet), the solution is to simply add __device__ qualifier for constructor:

__device__ cuComplex( float a, float b ) : r(a), i(b)  {}