CUDA By Examples 3 - 绘制Julia Set

1. 使用CPU方法

CUDA By Examples 第四章例程:

#include "../common/book.h"
#include "../common/cpu_bitmap.h"
//不能超过65535
#define DIM 512

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

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;
}

void kernel( unsigned char *ptr ){
for (int y=0; y<DIM; y++) {
for (int x=0; x<DIM; x++) {
int offset = x + y * DIM;

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;
}
}
}

int main( void ) {
CPUBitmap bitmap( DIM, DIM );
unsigned char *ptr = bitmap.get_ptr();

kernel( ptr );

bitmap.display_and_exit();
}

还需要添加随书代码包里的头文件和库文件等.

编译时会和MSVCRTD.LIB冲突, 需要在工程配置中忽略此库:



运行结果:

2. 使用GPU方法

__device__

函数在device上运行, 不能用host调用, 也不能调用host的函数.

__global__

函数在device上运行, 由host调用.

#include "../common/book.h"
#include "../common/cpu_bitmap.h"

#define DIM 512

struct cuComplex {
float   r;
float   i;
//改正: cuComplex必须是__device__, 因为被__device__ julia 调用了
__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();
}

输出和CPU方法一样.

补充: CUDA函数前缀详解

1.

__device__

使用

_device_

限定符声明的函数具有以下特征：

a. 在设备上执行；

b.仅可通过设备调用。

2.

__global__

使用

_global_

限定符可将函数声明为内核。此类函数：

a. 在设备上执行；

b. 仅可通过主机调用。

3.

__host__

使用

_host_

限定符声明的函数具有以下特征：

a. 在主机上执行；

b. 仅可通过主机调用。