cuda by example -julia

    在学习cuda by example 第四章julia例子过程中遇到的问题：

• cuComplex（float a, float b): r(a),r(b) {}，为C++类构造函数初始化列表。构造函数初始化列表以一个冒号开始，接着是以逗号分隔的数据成员列表，每个数据成员后面跟一个放在括号中的初始化式。C++初始化类成员时，是按照声明的顺序初始化的，而不是按照出现在初始化列表中的顺序。
• 对于在GPU中构造函数的初始化需要在前面加上__device__ 声明,至于__device__和__global__区别如下：__device__标记的函数从一个在器件中执行的函数呼叫，在器件中执行；__global__标示该函数从一个在主机中执行的函数呼叫，在器件中执行，例如__global__ void add(int a, int b, int *c) {*c=a+b;}是在CPU端的代码，声明add在GPU端执行； __host__标示在主机中呼叫，在主机中执行的函数。

      

#include "cuda_runtime.h"

#include "device_launch_parameters.h"
#include "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.5\common\book.h"
#include "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.5\common\cpu_bitmap.h"
#include <stdio.h>

#define DIM 1000

struct cuComplex {
float r;
float i;
__device__ 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();
}