纹理内存是只读内存，与常量内存相同的是，纹理内存也缓存在芯片中，因此某些情况下，它能减少对内存的请求并提供更高效的内存宽带。纹理内存专门为那些内存访问模式中存在大量空间局部性的图形应用程序而设计的。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。纹理缓存为了加速访问不连续的地址而设计的。
首先是一个不使用纹理内存的热传导模型。先用到了下面的辅助函数：
#ifndef __ANIM_H__
#define __ANIM_H__
#include <iostream>
#include <windows.h>
#include <GL/glut.h>
struct AnimBitmap {
unsigned char *pixels; //pixels是像素数组，大小为w*h*4
int width, height;
void *dataBlock; //需要处理的数据块
void (*fAnim)(void*,int); //函数指针fAnim，在glutIdleFunc注册后，空闲时一直执行它（动画）
void (*animExit)(void*); //函数指针animExit，在响应退出的时候可以调用这个函数
void (*clickDrag)(void*,int,int,int,int); //函数指针clickDrag，响应拖动的处理
int dragStartX, dragStartY;
AnimBitmap( int w, int h, void *d = NULL ) { //构造函数，几个函数指针暂时为NULL
width = w;
height = h;
pixels = new unsigned char[width * height * 4];
dataBlock = d;
clickDrag = NULL;
}
~AnimBitmap() {
delete [] pixels;
}
unsigned char* get_ptr( void ) const { return pixels; }
long image_size( void ) const { return width * height * 4; }
void click_drag( void (*f)(void*,int,int,int,int)) { //设置拖动指针的值
clickDrag = f;
}
void anim_and_exit( void (*f)(void*,int), void(*e)(void*) ) {
AnimBitmap** bitmap = get_bitmap_ptr();
*bitmap = this; //让静态变量所指的内容为当前定义的Bitmap
fAnim = f; //设置动画函数
animExit = e; //设置退出需要处理的函数
int c=1;
char* dummy = "";
glutInit( &c, &dummy );
glutInitDisplayMode( GLUT_DOUBLE | GLUT_RGBA );
glutInitWindowSize( width, height );
glutCreateWindow( "bitmap" );
glutKeyboardFunc(Key); //注册按键响应
glutDisplayFunc(Draw); //注册画面内容
if (clickDrag != NULL)
glutMouseFunc( mouse_func ); //注册拖动处理
glutIdleFunc( idle_func );//当没有窗口事件到达时，glutIdleFunc可以执行后台处理任务或连续动画。它被不断调用，直到有窗口事件发生。
glutMainLoop();
}
static AnimBitmap** get_bitmap_ptr( void ) {
static AnimBitmap* gBitmap; //静态变量，只第一次执行
return &gBitmap;
}
static void mouse_func( int button, int state, int mx, int my ) {
if (button == GLUT_LEFT_BUTTON) {
AnimBitmap* bitmap = *(get_bitmap_ptr());
if (state == GLUT_DOWN) {
bitmap->dragStartX = mx;
bitmap->dragStartY = my;
} else if (state == GLUT_UP) {
bitmap->clickDrag( bitmap->dataBlock,bitmap->dragStartX,bitmap->dragStartY,mx, my );
}
}
}
static void idle_func( void ) {
static int ticks = 1;
AnimBitmap* bitmap = *(get_bitmap_ptr());
bitmap->fAnim( bitmap->dataBlock, ticks++ );
glutPostRedisplay();
}
static void Key(unsigned char key, int x, int y) {
switch (key) {
case 27:
AnimBitmap* bitmap = *(get_bitmap_ptr());
bitmap->animExit( bitmap->dataBlock );
//delete bitmap;
exit(0);
}
}
static void Draw( void ) {
AnimBitmap* bitmap = *(get_bitmap_ptr());
glClearColor( 0.0, 0.0, 0.0, 1.0 );
glClear( GL_COLOR_BUFFER_BIT );
glDrawPixels( bitmap->width, bitmap->height, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
glutSwapBuffers();
}
};
#endif
这个AnimBitmap.h文件用位图的指针来生成图，并产生动画。
例子中还是用了一个辅助函数，这个函数是为了将浮点数转化成RGBA值：
__device__ unsigned char value( float n1, float n2, int hue ) {
if (hue > 360) hue -= 360;
else if (hue < 0) hue += 360;
if (hue < 60)
return (unsigned char)(255 * (n1 + (n2-n1)*hue/60));
if (hue < 180)
return (unsigned char)(255 * n2);
if (hue < 240)
return (unsigned char)(255 * (n1 + (n2-n1)*(240-hue)/60));
return (unsigned char)(255 * n1);
}
__global__ void float_to_color( unsigned char *optr,const float *outSrc ) {
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 l = outSrc[offset];
float s = 1;
int h = (180 + (int)(360.0f * outSrc[offset])) % 360;
float m1, m2;
if (l <= 0.5f)
m2 = l * (1 + s);
else
m2 = l + s - l * s;
m1 = 2 * l - m2;
optr[offset*4 + 0] = value( m1, m2, h+120 );
optr[offset*4 + 1] = value( m1, m2, h );
optr[offset*4 + 2] = value( m1, m2, h -120 );
optr[offset*4 + 3] = 255;
}
需要处理的数据块是BlockData，这个数据块是需要处理的，通过这些数据得到最后位图，在退出时候，进行回收。
#include "AnimBitmap.h"
#define DIM 1024
#define SPEED 0.25
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
__global__ void copy_const_kernel(float *iptr,const float *cptr){
int x=threadIdx.x+blockIdx.x*blockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;
int offset=x+y*blockDim.x*gridDim.x; //输入缓冲区的线性偏移
if(cptr[offset]!=0)
iptr[offset]=cptr[offset]; //当温度非零时，将温度复制到iptr中
}
__global__ void blend_kernel(float *outSrc,const float *inSrc){
int x=threadIdx.x+blockIdx.x*blockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;
int offset=x+y*blockDim.x*gridDim.x;
//每个线程与左右上下四点对应
int left=offset-1;
int right=offset+1;
if(x==0) left++;
if(x==DIM-1) right--;
int top=offset-DIM;
int bottom=offset+DIM;
if(y==0) top+=DIM;
if(y==DIM-1) bottom-=DIM;
//温度传播
outSrc[offset]=inSrc[offset]+SPEED*(inSrc[top]+inSrc[bottom]+inSrc[left]+inSrc[right]-inSrc[offset]*4);
}
struct DataBlock{
unsigned char* output_bitmap;
float *dev_inSrc;
float *dev_outSrc;
float *dev_constSrc;
AnimBitmap *bitmap;
cudaEvent_t start,stop;
float totalTime;
float frames;
};
void anim(DataBlock *d,int ticks){
cudaEventRecord(d->start,0);
dim3 blocks(DIM/16,DIM/16);
dim3 threads(16,16);
AnimBitmap *bitmap=d->bitmap;
for(int i=0;i<90;i++){
copy_const_kernel<<<blocks,threads>>>(d->dev_inSrc,d->dev_constSrc);
blend_kernel<<<blocks,threads>>>(d->dev_outSrc,d->dev_inSrc);
float * current=d->dev_inSrc;
d->dev_inSrc=d->dev_outSrc;
d->dev_outSrc=current;
}
float_to_color<<<blocks,threads>>>(d->output_bitmap,d->dev_inSrc);
cudaMemcpy(bitmap->get_ptr(),d->output_bitmap,bitmap->image_size(),cudaMemcpyDeviceToHost);
cudaEventRecord(d->stop,0);
cudaEventSynchronize(d->stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,d->start,d->stop);
d->totalTime+=elapsedTime;
++d->frames;
printf("Average Time per frame:%3.1f ms\n",d->totalTime/d->frames);
}
void anim_exit(DataBlock *d){
cudaFree(d->dev_inSrc);
cudaFree(d->dev_outSrc);
cudaFree(d->dev_constSrc);
cudaEventDestroy(d->start);
cudaEventDestroy(d->stop);
}
int main(void){
DataBlock data;
AnimBitmap bitmap(DIM,DIM,&data);
data.bitmap=&bitmap;
data.totalTime=0;
data.frames=0;
cudaEventCreate(&data.start);
cudaEventCreate(&data.stop);
cudaMalloc((void**)&data.output_bitmap,bitmap.image_size());
cudaMalloc((void**)&data.dev_inSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_outSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_constSrc,bitmap.image_size());
float *temp=(float*)malloc(bitmap.image_size());
for(int i=0;i<DIM*DIM;i++){
temp[i]=0;
int x=i%DIM;
int y=i/DIM;
if((x>300)&&(x<600)&&(y>310)&&(y<601))
temp[i]=MAX_TEMP;
}
temp[DIM*100+100]=(MAX_TEMP+MIN_TEMP)/2;
temp[DIM*700+100]=MIN_TEMP;
temp[DIM*300+300]=MIN_TEMP;
temp[DIM*200+700]=MIN_TEMP;
for(int y=800;y<900;y++)
for(int x=400;x<500;x++)
temp[x+y*DIM]=MIN_TEMP;
cudaMemcpy(data.dev_constSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice); //dev_constSrc是热源
for(int y=800;y<DIM;y++)
for(int x=0;x<200;x++)
temp[x+y*DIM]=MAX_TEMP;
cudaMemcpy(data.dev_inSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice);//设置初始温度
free(temp);
bitmap.anim_and_exit( (void (*)(void*,int))anim,(void (*)(void *))anim_exit );
return 0;
}
接下来使用纹理内存：
温度计算的内存访问模式中存在着巨大的内存空间局部性，这种访问模式可以用GPU纹理内存加速。首先声明纹理内存 texture<float> tex；这个缓存区域分配内存后需要绑定到内存缓冲区。然后，启动核函数时，要用特殊的函数告诉GPU将读取请求转发到纹理内存而不是标准全局内存。当读取内存时不再使用方括号冲缓冲区中读取，而是将blend_kernel()改为tex1Dfetch().
blend_kernel()中又一个参数dstOut告诉那个缓冲区作为输入，哪个作为输出。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include "AnimBitmap.h"
#define DIM 1024
#define SPEED 0.25f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
texture<float> texConst;
texture<float> texIn;
texture<float> texOut;
__global__ void copy_const_kernel(float *iptr){
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 c=tex1Dfetch(texConst,offset);
if(c>0)
iptr[offset]=c;
}
__global__ void blend_kernel(float *dst,bool dstOut){
int x=threadIdx.x+blockIdx.x*blockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;
int offset=x+y*blockDim.x*gridDim.x;
int left=offset-1;
int right=offset+1;
if(x==0) left++;
if(x==DIM-1) right--;
int top=offset-DIM;
int bottom=offset+DIM;
if(y==0) top+=DIM;
if(y==DIM-1) bottom-=DIM;
float t,l,c,r,b;
if(dstOut){ //true时候，texIn的区域作为输入
t=tex1Dfetch(texIn,top);
l=tex1Dfetch(texIn,left);
c=tex1Dfetch(texIn,offset);
r=tex1Dfetch(texIn,right);
b=tex1Dfetch(texIn,bottom);
}else{ //false时候，texOut的区域作为输入
t=tex1Dfetch(texOut,top);
l=tex1Dfetch(texOut,left);
c=tex1Dfetch(texOut,offset);
r=tex1Dfetch(texOut,right);
b=tex1Dfetch(texOut,bottom);
}
dst[offset]=c+SPEED*(t+b+r+l-4*c);
}
struct DataBlock{
unsigned char* output_bitmap;
float *dev_inSrc;
float *dev_outSrc;
float *dev_constSrc;
AnimBitmap *bitmap;
cudaEvent_t start,stop;
float totalTime;
float frames;
};
__device__ unsigned char value( float n1, float n2, int hue ) {
if (hue > 360) hue -= 360;
else if (hue < 0) hue += 360;
if (hue < 60)
return (unsigned char)(255 * (n1 + (n2-n1)*hue/60));
if (hue < 180)
return (unsigned char)(255 * n2);
if (hue < 240)
return (unsigned char)(255 * (n1 + (n2-n1)*(240-hue)/60));
return (unsigned char)(255 * n1);
}
__global__ void float_to_color( unsigned char *optr,const float *outSrc ) {
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 l = outSrc[offset];
float s = 1;
int h = (180 + (int)(360.0f * outSrc[offset])) % 360;
float m1, m2;
if (l <= 0.5f)
m2 = l * (1 + s);
else
m2 = l + s - l * s;
m1 = 2 * l - m2;
optr[offset*4 + 0] = value( m1, m2, h+120 );
optr[offset*4 + 1] = value( m1, m2, h );
optr[offset*4 + 2] = value( m1, m2, h -120 );
optr[offset*4 + 3] = 255;
}
void anim(DataBlock *d,int ticks){
cudaEventRecord(d->start,0);
dim3 blocks(DIM/16,DIM/16);
dim3 threads(16,16);
AnimBitmap *bitmap=d->bitmap;
volatile bool dstOut=true;
float *in,*out;
for(int i=0;i<90;i++){
if(dstOut){
in=d->dev_inSrc;
out=d->dev_outSrc;
}else{
out=d->dev_inSrc;
in=d->dev_outSrc;
}
copy_const_kernel<<<blocks,threads>>>(in);
blend_kernel<<<blocks,threads>>>(out,dstOut);
dstOut=!dstOut;
}
float_to_color<<<blocks,threads>>>(d->output_bitmap,d->dev_inSrc);
cudaMemcpy(bitmap->get_ptr(),d->output_bitmap,bitmap->image_size(),cudaMemcpyDeviceToHost);
cudaEventRecord(d->stop,0);
cudaEventSynchronize(d->stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime,d->start,d->stop);
d->totalTime+=elapsedTime;
++d->frames;
printf("Average Time per frame:%3.1f ms\n",d->totalTime/d->frames);
}
void anim_exit(DataBlock *d){
//清楚绑定
cudaUnbindTexture(texIn);
cudaUnbindTexture(texOut);
cudaUnbindTexture(texConst);
cudaFree(d->dev_inSrc);
cudaFree(d->dev_outSrc);
cudaFree(d->dev_constSrc);
cudaEventDestroy(d->start);
cudaEventDestroy(d->stop);
}
int main(void){
DataBlock data;
AnimBitmap bitmap(DIM,DIM,&data);
data.bitmap=&bitmap;
data.totalTime=0;
data.frames=0;
cudaEventCreate(&data.start);
cudaEventCreate(&data.stop);
cudaMalloc((void**)&data.output_bitmap,bitmap.image_size());
cudaMalloc((void**)&data.dev_inSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_outSrc,bitmap.image_size());
cudaMalloc((void**)&data.dev_constSrc,bitmap.image_size());
//绑定纹理变量到内存的缓冲区
cudaBindTexture(NULL,texConst,data.dev_constSrc,bitmap.image_size());
cudaBindTexture(NULL,texIn,data.dev_inSrc,bitmap.image_size());
cudaBindTexture(NULL,texOut,data.dev_outSrc,bitmap.image_size());
float *temp=(float*)malloc(bitmap.image_size());
for(int i=0;i<DIM*DIM;i++){
temp[i]=0;
int x=i%DIM;
int y=i/DIM;
if((x>300)&&(x<600)&&(y>310)&&(y<601))
temp[i]=MAX_TEMP;
}
cudaMemcpy(data.dev_constSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice);
cudaMemcpy(data.dev_inSrc,temp,bitmap.image_size(),cudaMemcpyHostToDevice);
free(temp);
bitmap.anim_and_exit( (void (*)(void*,int))anim,(void (*)(void *))anim_exit );
return 0;
}
除了一维纹理，还有二维纹理。使用是类似的。
定义的时候用texture<float,2> tex;同样进行绑定cudaBindTexture2D(NULL,tex,dev_inSrc,desc,DIM,DIM,size)
最后解除绑定。在写核函数时候，这里有一个好处是不用处理边界。
#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED 0.25f
// these exist on the GPU side
texture<float,2> texConstSrc;
texture<float,2> texIn;
texture<float,2> texOut;
__global__ void blend_kernel( float *dst,
bool dstOut ) {
// map from threadIdx/BlockIdx to pixel position
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 t, l, c, r, b;
if (dstOut) {
t = tex2D(texIn,x,y-1);
l = tex2D(texIn,x-1,y);
c = tex2D(texIn,x,y);
r = tex2D(texIn,x+1,y);
b = tex2D(texIn,x,y+1);
} else {
t = tex2D(texOut,x,y-1);
l = tex2D(texOut,x-1,y);
c = tex2D(texOut,x,y);
r = tex2D(texOut,x+1,y);
b = tex2D(texOut,x,y+1);
}
dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}
__global__ void copy_const_kernel( float *iptr ) {
// map from threadIdx/BlockIdx to pixel position
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 c = tex2D(texConstSrc,x,y);
if (c != 0)
iptr[offset] = c;
}
// globals needed by the update routine
struct DataBlock {
unsigned char *output_bitmap;
float *dev_inSrc;
float *dev_outSrc;
float *dev_constSrc;
CPUAnimBitmap *bitmap;
cudaEvent_t start, stop;
float totalTime;
float frames;
};
void anim_gpu( DataBlock *d, int ticks ) {
HANDLE_ERROR( cudaEventRecord( d->start, 0 ) );
dim3 blocks(DIM/16,DIM/16);
dim3 threads(16,16);
CPUAnimBitmap *bitmap = d->bitmap;
// since tex is global and bound, we have to use a flag to
// select which is in/out per iteration
volatile bool dstOut = true;
for (int i=0; i<90; i++) {
float *in, *out;
if (dstOut) {
in = d->dev_inSrc;
out = d->dev_outSrc;
} else {
out = d->dev_inSrc;
in = d->dev_outSrc;
}
copy_const_kernel<<<blocks,threads>>>( in );
blend_kernel<<<blocks,threads>>>( out, dstOut );
dstOut = !dstOut;
}
float_to_color<<<blocks,threads>>>( d->output_bitmap,
d->dev_inSrc );
HANDLE_ERROR( cudaMemcpy( bitmap->get_ptr(),
d->output_bitmap,
bitmap->image_size(),
cudaMemcpyDeviceToHost ) );
HANDLE_ERROR( cudaEventRecord( d->stop, 0 ) );
HANDLE_ERROR( cudaEventSynchronize( d->stop ) );
float elapsedTime;
HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
d->start, d->stop ) );
d->totalTime += elapsedTime;
++d->frames;
printf( "Average Time per frame: %3.1f ms\n",
d->totalTime/d->frames );
}
// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
cudaUnbindTexture( texIn );
cudaUnbindTexture( texOut );
cudaUnbindTexture( texConstSrc );
HANDLE_ERROR( cudaFree( d->dev_inSrc ) );
HANDLE_ERROR( cudaFree( d->dev_outSrc ) );
HANDLE_ERROR( cudaFree( d->dev_constSrc ) );
HANDLE_ERROR( cudaEventDestroy( d->start ) );
HANDLE_ERROR( cudaEventDestroy( d->stop ) );
}
int main( void ) {
DataBlock data;
CPUAnimBitmap bitmap( DIM, DIM, &data );
data.bitmap = &bitmap;
data.totalTime = 0;
data.frames = 0;
HANDLE_ERROR( cudaEventCreate( &data.start ) );
HANDLE_ERROR( cudaEventCreate( &data.stop ) );
int imageSize = bitmap.image_size();
HANDLE_ERROR( cudaMalloc( (void**)&data.output_bitmap,
imageSize ) );
// assume float == 4 chars in size (ie rgba)
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,
imageSize ) );
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,
imageSize ) );
HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,
imageSize ) );
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
HANDLE_ERROR( cudaBindTexture2D( NULL, texConstSrc,
data.dev_constSrc,
desc, DIM, DIM,
sizeof(float) * DIM ) );
HANDLE_ERROR( cudaBindTexture2D( NULL, texIn,
data.dev_inSrc,
desc, DIM, DIM,
sizeof(float) * DIM ) );
HANDLE_ERROR( cudaBindTexture2D( NULL, texOut,
data.dev_outSrc,
desc, DIM, DIM,
sizeof(float) * DIM ) );
// initialize the constant data
float *temp = (float*)malloc( imageSize );
for (int i=0; i<DIM*DIM; i++) {
temp[i] = 0;
int x = i % DIM;
int y = i / DIM;
if ((x>300) && (x<600) && (y>310) && (y<601))
temp[i] = MAX_TEMP;
}
temp[DIM*100+100] = (MAX_TEMP + MIN_TEMP)/2;
temp[DIM*700+100] = MIN_TEMP;
temp[DIM*300+300] = MIN_TEMP;
temp[DIM*200+700] = MIN_TEMP;
for (int y=800; y<900; y++) {
for (int x=400; x<500; x++) {
temp[x+y*DIM] = MIN_TEMP;
}
}
HANDLE_ERROR( cudaMemcpy( data.dev_constSrc, temp,
imageSize,
cudaMemcpyHostToDevice ) );
// initialize the input data
for (int y=800; y<DIM; y++) {
for (int x=0; x<200; x++) {
temp[x+y*DIM] = MAX_TEMP;
}
}
HANDLE_ERROR( cudaMemcpy( data.dev_inSrc, temp,
imageSize,
cudaMemcpyHostToDevice ) );
free( temp );
bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,
(void (*)(void*))anim_exit );
}
相关推荐
CUDA by Example (2010).pdf CUDA by Example (2010).pdf
《CUDA By Example》中文译名《GPU高性能编程CUDA实战》是研究GPGPU异构并行计算非常不错的工具书。书中给出的代码，非常个别的地方有失误，但是都有人为标注了，而且对不同的编程工具可能需要自己配置链接库。...
CUDA by example代码实例
cuda by example 中英文版及代码，中文版名称是 GPU高性能编程CUDA实战，含书签
《CUDA By Example》中文译名《GPU高性能编程CUDA实战》 源码 包括 book.h cpu_anim.h cpu_bitmap.h gl_helper.h gpu_anim.h glext.h glut.h
通过实例学习cuda编程，nvidia GPU学习教程...........
本资源提供了CUDA By Example一书中用到的源代码
cuda by example,中文版，是好东西哦~
CUDA by Example.An Introduction to General-Purpose GPU Programming（英文原书+自带源代码）源代码是nvidia官网下的。
CUDA by Example: An Introduction to General-Purpose GPU Programming
cuda 经典的书 很好的入门教程 建议学习gpu的可以看看
《GPU高性能编程 CUDA实战》/《CUDA By Example》中的book.h
CUDA By Example书中的有book.H会无处找，以上就是书中为给出的代码。该代码是NVIDIA的代码
CUDA by Example.An Introduction to General-Purpose GPU Programming 需要的自然懂 pdf 版， 带目录， 全书313页~
CUDA by Example: An Introduction to General-Purpose GPU Programming带随书源代码。
《CUDA BY EXAMPLE》（GPU高性能编程CUDA实战）书中所有EXAMPLE的代码，包括书中提供的库。
《GPU高性能编程CUDA实战（CUDA By Example）》头文件_全书使用的，在163页通过多线程，使用多GPU的地方已完美调试，可运行通过。
介绍GPU编程，给出了很多CUDA编程的示例