在GPU上使用线程实现波纹效果

#include "cuda.h"
#include "book.h"
#include "cpu_anim.h"

#define DIM 1024
#define PI 3.1415926

struct DataBlock
{
 unsigned char *dev_bitmap;
 CPUAnimBitmap *bitmap;
};
//释放在GPU上分配的内存
void cleanup(DataBlock *d)
{
 cudaFree(d->dev_bitmap);
}

__global__ void kernel(unsigned char *ptr,int ticks)
{
 //将threadIdx/BlockIdx映射到像素位置
 int x=threadIdx.x+blockIdx.x*blockDim.x;
 int y=threadIdx.y+blockIdx.y*blockDim.y;
 int offset=x+y*blockDim.x*gridDim.x;
 //计算当前位置上的值
 float fx=x-DIM/2;
 float fy=y-DIM/2;
 float d=sqrt(fx*fx+fy*fy);
 unsigned char grey=(unsigned char)(128.0f+127.0f*cos(d/10.0f-ticks/7.0f)/(d/10.0f+1.0f));
 ptr[offset*4+0]=grey;
 ptr[offset*4+1]=grey;
 ptr[offset*4+2]=grey;
 ptr[offset*4+3]=255;
}
void generate_frame(DataBlock *d,int ticks)
{
 dim3 blocks(DIM/16,DIM/16);
 dim3 threads(16,16);
 kernel<<<blocks,threads>>>(d->dev_bitmap,ticks);
 HANDLE_ERROR(cudaMemcpy(d->bitmap->get_ptr(),d->dev_bitmap,d->bitmap->image_size(),cudaMemcpyDeviceToHost));
}
int main(void)
{
 DataBlock data;
 CPUAnimBitmap bitmap(DIM,DIM,&data);
 data.bitmap=&bitmap;
 HANDLE_ERROR(cudaMalloc((void**)&data.dev_bitmap,bitmap.image_size()));
 bitmap.anim_and_exit((void(*)(void*,int))generate_frame,(void(*)(void*))cleanup);
}