4.2 CUDA并行编程
4.2.1 矢量求和运算
- 基于CPU的矢量求和
代码:add_loop_cpu.cu - 基于GPU的矢量求和
代码:add_loop_gpu.cu
add<<<N,1>>>(dev_a, dev_b, dev_c);
尖括号中第一个数值:设备在执行核函数时使用的并行线程块的数量。
blockIdx 是一个内置变量,CUDA运行时已经预先定义了这个变量。
一维的线程格,其中包含N维的线程块。
4.2.2 一个有趣的示例
__global__ 函数 从主机上调用并在设备上运行。
__device__ 函数 在设备上运行。从其他 __device__函数或者__global__函数调用它们。
代码:
#include "../common/book.h"
#include "../common/cpu_bitmap.h"
#define DIM 1000
struct cuComplex {
float r;
float i;
// cuComplex( float a, float b ) : r(a), i(b) {}
__device__ cuComplex( float a, float b ) : r(a), i(b) {} // Fix error for calling host function from device
__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;
printf("blockIdx.x %d, blockIdx.y %d \n", x, 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 );
// 指定了多个并行线程块来执行函数kernel()。
HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
bitmap.image_size(),
cudaMemcpyDeviceToHost ) );
HANDLE_ERROR( cudaFree( dev_bitmap ) );
bitmap.display_and_exit();
}