在运行《CUDA By Example》中的julia_gpu.cu时出现CUDA错误。

Question

我正在运行书籍《CUDA By Example》中的示例julia_gpu.cu，使用的是Visual Studio Express 2012中的CUDA 6.0。以下是参考源代码：

#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();
}

编译上述内容时出现了一系列错误，这些错误都具有相同的形式：

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

我很难理解问题所在，因为据我所知，__device__函数（例如julia）可以自由调用和创建struct对象（例如cuComplex）。这里到底是什么问题？《CUDA By Example》中提供的代码（写于2010年）是否被更近期的CUDA 6.0更新中引入的一些更改破坏了？

- DumpsterDoofus

Answer 1

我之前已经向CUDA团队提交了勘误（但尚未收到任何回复），解决方案是简单地为构造函数添加__device__限定符：

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

- Grzegorz Szpetkowski

---

那么这是否意味着在编写该书时，不需要__device__限定符，但在后来的CUDA版本中变得必要了呢？还是只是一个打字错误？感谢您的回复，它起作用了。 - DumpsterDoofus

@DumpsterDoofus：这只是一个打字错误。 - Grzegorz Szpetkowski

Answer 2

您正在从设备函数调用cuComplex构造函数，但这是不可能的，因为它是一个主机函数。

您可以尝试将构造函数定义为设备函数，但我不确定它是否会起作用。

- MoonBun