cudaMalloc是在设备上分配内存，第一个参数是新分配内存的地址，第二个参数是分配内存大小。在主机上不能对这块内存做任何的修改。主机指针只能访问主机代码中的内存，设备指针只能访问设备代码中的内存。使用完设备内存后要用cudaFree来释放掉分配的内存。在主机代码中可以调用cudaMemcpy来访问设备上的内存。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<conio.h>
#include <stdio.h>
__global__ void add(int a,int b,int *c){ //设备上运行的函数
*c=a+b;
}
int main(){
int c;
int *dev_c;
cudaMalloc((void**)&dev_c,sizeof(int)); //在设备上给指针dev_c所指向单元分配内存
add<<<1,1>>>(2,7,dev_c); //一个grid中只有一个block，一个block中只有一个thread模式来运行kernel
cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost); //设备上dev_c所指内容赋给主机上的c
printf("2+7=%d",c);
cudaFree(dev_c); //释放设备dev_c所指的内存
getche();
return 0;
}
尖括号中有两个参数，第一个参数是指设备执行核函数时候使用的并行线程块block的数量。kernel<<<256,1>>>意味着一个Grid中将有256个线程块来执行这个kernel。可以用blockIdx.x来获取线程块的索引。CUDA支持二维线程块数组，在处理矩阵时候或者图形的时候是非常便利的。在启动线程块数组时，数组的每一维数不能超过65335.第二个参数表示运行在每个线程块中创建的线程数量。线程块中的并行线程能够完成并行线程块（block）无法完成的工作。线程块中的线程可以用threadIdx.x来索引。
#ifndef _use_h
#define _use_h
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <conio.h>
#include <stdio.h>
#define N 100
__global__ void add(int *a,int *b,int *c);
#endif
#include "use.h"
__global__ void add(int *a,int *b,int *c){
//int tid=blockIdx.x; //tid为第i个block
int tid=threadIdx.x; //tid为第i个thread
if(tid<N)
c[tid]=a[tid]+b[tid];
}
#include "use.h"
int main(){
int a[N],b[N],c[N];
int *dev_a,*dev_b,*dev_c;
cudaMalloc((void**)&dev_a,N*sizeof(int));
cudaMalloc((void**)&dev_b,N*sizeof(int));
cudaMalloc((void**)&dev_c,N*sizeof(int));
//CPU上a,b赋值
for(int i=0;i<N;i++){
a[i]=-i;
b[i]=i*i;
}
//将a，b拷贝到GPU上
cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);
//add<<<N,1>>>(dev_a,dev_b,dev_c); //N个block，每个block中一个线程的模式运行add
add<<<1,N>>>(dev_a,dev_b,dev_c); //1个block，每个block中有N个线程的模式运行add
//将c从GPU复制到CPU中
cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);
for(int i=0;i<N;i++){
printf("%d+%d=%d\n",a[i],b[i],c[i]);
}
//释放GPU上分配的空间
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
getche();
}
可以用gridDim.x获取每个grid中有多少block，用blockDim.x来获取每个block中有多少thread。这两个数在线程定位中很有用。以上向量相加并没有考虑线程数量与硬件的一些限制，例如线程格每一位不超过65535。以下对其改进：
#include "use.h"
__global__ void add(int *a,int *b,int *c){
int tid=blockIdx.x*blockDim.x+threadIdx.x; //第i个block前面有i*blockDim个线程，tid定位到当前的线程
while(tid<N){
c[tid]=a[tid]+b[tid];
tid+=blockDim.x*gridDim.x; //如果N相对线程数任然很大，那么可以让第i个线程处理第（i+k*线程数）个线程
}
}
#include "use.h"
int main(){
int a[N],b[N],c[N];
int *dev_a,*dev_b,*dev_c;
//GPU上分配内存
cudaMalloc((void**)&dev_a,N*sizeof(int));
cudaMalloc((void**)&dev_b,N*sizeof(int));
cudaMalloc((void**)&dev_c,N*sizeof(int));
//CPU上a,b赋值
for(int i=0;i<N;i++){
a[i]=i;
b[i]=i*i;
}
//将a，b拷贝到GPU上
cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);
add<<<128,128>>>(dev_a,dev_b,dev_c); //128个block，每个block中有128个线程的模式运行add，共16384，约为N的1/3
//将c从GPU复制到CPU中
cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);
int success=1;
for(int i=0;i<N;i++){
//printf("%d+%d=%d\n",a[i],b[i],c[i]);
if(a[i]+b[i]!=c[i])
{success=0;break;}
}
if(success==1)
printf("sucess\n");
else
printf("failure\n");
//释放GPU上分配的空间
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
getche();
}
关键字__share__用于变量声明，可以将变量驻留在共享内存中。CUDA C对共享变量与普通变量不同处理。对于共享内存中的变量，编译器会创建它的一个副本，这个block中的每个thread共享这个副本（别的block中的线程不能访问），这样共享内存缓冲区驻留在物理GPU中。访问共享内存的延迟比访问普通缓冲区的延迟要远远低。但线程之间的通信还要实现同步机制。
同步机制用__syncthreads()来实现。代码中如果使用了这个函数，那么，只有每个线程执行了这个函数以后，才允许进一步执行其余的部分。
下面这个例子是计算两个向量的内积：
#ifndef _use_h
#define _use_h
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <conio.h>
#include <stdio.h>
#include <stdlib.h>
#define imin(a,b) (a<b?a:b)
const int N=33*1024;
const int threadPerBlock=256;
const int blockPerGrid=imin(32,(N+threadPerBlock-1) / threadPerBlock);
__global__ void dot(float *a,float *b,float *c);
#endif
#include "use.h"
__global__ void dot(float *a,float *b,float *c){
__shared__ float cache[threadPerBlock]; //共享变量，存储a[i]*b[i]的结果，每个block中都有这个变量
int tid=threadIdx.x+blockIdx.x*blockDim.x; //一个grid中的线程号码
int cacheIndex=threadIdx.x; //block中的线程编号
float temp=0;
while(tid<N){
temp+=a[tid]*b[tid]; //自己的线程处理的数的和相加
tid+=blockDim.x*gridDim.x; //tid加上所以线程数
}
cache[cacheIndex]=temp; //确定block中变量chace的值
__syncthreads(); //让线程块中的所以线程同步，即让所有线程执行到这个位置
int i=blockDim.x/2;
while(i!=0){
if(cacheIndex<i)
cache[cacheIndex]+=cache[cacheIndex+i];
__syncthreads(); //cache[0]到cache[i]的值都确定
i=i/2;
}
if(cacheIndex==0)
c[blockIdx.x]=cache[0];
}
#include "use.h"
int main(){
float *a,*b,c,*partial_c; //a,b是要求内积的向量，c是和，partial_c长为grid中的block数，
//partial_c[i]是是第i个block每个线程求得的数据的和
float *dev_a,*dev_b,*dev_partial_c;
//在CPU上分配内存
a=(float*)malloc(N*sizeof(float));
b=(float*)malloc(N*sizeof(float));
partial_c=(float*)malloc(blockPerGrid*sizeof(float));
//在GPU上分配内存
cudaMalloc((void**)&dev_a,N*sizeof(float));
cudaMalloc((void**)&dev_b,N*sizeof(float));
cudaMalloc((void**)&dev_partial_c,blockPerGrid*sizeof(float));
//填充主机内存
for(int i=0;i<N;i++){
a[i]=i;
b[i]=i*2;
}
//将a，b复制到GPU
cudaMemcpy(dev_a,a,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(dev_b,b,N*sizeof(float),cudaMemcpyHostToDevice);
dot<<<blockPerGrid,threadPerBlock>>>(dev_a,dev_b,dev_partial_c);
//复制partial_c到HOST上，在HOST上求和
cudaMemcpy(partial_c,dev_partial_c,blockPerGrid*sizeof(float),cudaMemcpyDeviceToHost);
c=0;
for(int i=0;i<blockPerGrid;i++){
c+=partial_c[i];
}
#define sum_squares(x) (x*(x+1)*(2*x+1)/6)
printf("Does GPU value %.6g=%.6g?\n",c,2*sum_squares((float)(N-1)));
//释放GPU上内存
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_partial_c);
free(a);
free(b);
free(partial_c);
getche();
}
这个求内积过程中抽象为：对于一个输入数组进行某种计算，然后产生一个更小的结果数组，这个过程也成为归约（reduction）。实现归约的最简单的办法是，由一个线程在共享内存上进行迭代并计算出总和值。
代码中;
int i=blockDim.x/2;
while(i!=0){
if(cacheIndex<i)
cache[cacheIndex]+=cache[cacheIndex+i];
__syncthreads();
i=i/2;
}
这里意味着只有部分线程进行工作，而其余的线程总是在进行__syncthreads等待，或许会有以下的优化，即让那些不做工作的线程继续执行至结束，而不是等待，这样或许效率高一些：
int i=blockDim.x/2;
while(i!=0){
if(cacheIndex<i){
cache[cacheIndex]+=cache[cacheIndex+i];
__syncthreads();
}
i=i/2;
}
但是，这是无法工作的，因为，__syncthreads必须保证每个线程执行到这里才会让线程进行下一步工作，但是，有些线程是永远都不会执行__syncthreads()，这样会让程序无尽地等待下去。
相关推荐
CUDA by Example (2010).pdf CUDA by Example (2010).pdf
本资源提供了CUDA By Example一书中用到的源代码
《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 编程实战 《cuda by example》一书中的头文件，是其自定义的。方便用户直接引用
《CUDA BY EXAMPLE》（GPU高性能编程CUDA实战）书中所有EXAMPLE的代码，包括书中提供的库。