CUDA绘制Julia集图像debug实录

症状：代码无法运行，报错显示为“请输入表达式”，代码原始状态如下

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cpu_bitmap.h"
#include "book.h"
 
#include <stdio.h>
 
#define DIM 1000
 
#define R1 214 //底色，根据自己喜好输入颜色RGB值
#define G1 163
#define B1 220
 
#define R2 247 //Julia集颜色，根据自己喜好输入颜色RGB值
#define G2 219
#define B2 112
 
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);
 
    for (int i = 0; i < 200; ++i) {
        a = a * a + c;
        //该点发散，不在Julia集中
        if (a.magnitude2() > 1000) return 0;
    }
    return 1;//该点在Julia集中
}
 
__global__ void Kernel(unsigned char* ptr)
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;
 
    int juliaValue = julia(x, y);
    //RGB1是底色，RGB2是Julia集中的点的颜色
    ptr[offset * 4 + 0] = R1 + (R2-R1)*juliaValue;//R
    ptr[offset * 4 + 1] = G1 + (G2-G1)*juliaValue;//G
    ptr[offset * 4 + 2] = B1 + (B2-B1)*juliaValue;//B
    //位图中每个像素的RGB保存在一个二维数组
    //为了凑整，使用4个字节存储一个RGB
    //第四个字节不使用，无效，可管理可不管理
    //ptr[offset * 4 + 3] = 255;
}
 
int main()
{
    CPUBitmap bitmap(DIM, DIM);
    unsigned char* dev_bitmap;
 
    HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, bitmap.image_size()));
 
    dim3 grid(DIM, DIM);
    Kernel << <grid, 1 >> > (dev_bitmap);
 
    HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,
        bitmap.image_size(),
        cudaMemcpyDeviceToHost));
 
    bitmap.display_and_exit();
 
    return 0;
}

提示错误出现在语句Kernel << <grid, 1 >> > (dev_bitmap);   上网查询，有说是因为编译器不同的原因，.cu文件和.cpp文件的编译是不同的。在.cpp文件与.cu文件混合编写时，不应该在.cpp文件里出现.cu文件的函数，而是应该使用外部拓展c标识符，而我的代码文件后缀名就是cu，所以说不是这个问题。出现这个报错还有个可能原因：__global__定义的核函数kernel中有未定义的标识符，结果这个error没有能够定位好，就让<<< >>>给背锅了！！我非常仔细的在核函数Kernel中检查各个标识符，但是并没有找到未定义的标识符。太奇怪了，到底是什么问题呢？我搜到一个可以正确运行的代码，一对比，才知道问题原来出在cuComplex结构体的构造函数！在构造函数前加__device__程序就正常运行了！这怎么解释呢？这和我之前查到的可能原因有关吗？

答案是有关。

因为cpu和gpu的内存不互通，所以分配到cpu内存的标识符gpu不能用，分配到gpu内存的标识符cpu也不能用。cuComplex的构造函数前没加__device__，所以构造是在cpu上运行的，构造出来的对象是属于cpu内存的，而julia函数要在gpu中运行，cpu构造出来的cuComplex对象gpu无法使用，julia函数也就无法使用，Kernel核函数又会调用julia函数，所以就相当于核函数Kernel中出现了未定义的标识符（未在gpu中定义），又因为error定位得不准确，<<<>>>就背锅了！