__global__ void VecAdd(float* A,float* B,float* C)
{}
int main()
{
VecAdd<<<1,N>>>(A,B,C);
}
其中，1，N是函数执行参数，代表kernel的Grid中只有一个block，而每个block中则有N个thread，而A,B，C是函数的参数。其中threadID和blockID都是内建的变量，用于和其他的线程相区分。Kernel是以block为单位执行的。
下面演示对两个N*N的矩阵A和B进行求和。并将结果保存在C中。
__global__ void MatAdd(float A[N][N],float B[N][N],float C[N][N])
{
int i=threadIdx.x;
int j=threadIdy.y;
c[i][j]=A[i][j]+B[i][j];
}
int main()
{
dim3 dimBlock(N,N);
MatAdd<<<1,dimBlock>>>(A,B,C);
}
这段代码说明使用多维的thread ID不仅可以使对多维的并行处理的编程更加方便，同时也可以避免一些求商和求余的工作。
cuda中实现block通信的方法是：在同一个block中的线程通过共享存储器交换数据，并通过栅栏同步保证线程间能够正确的共享数据。可以在kernel函数中需要同步的位置调用__syncthread（）函数。
__global__ void Sum(float* A,float* B)
{
int bid=blockIdx.x;
int tid=threadIdx.x;
__shared__ s_data[128];
//将数据从缓存读入shared memory
s_data[tid]=A[bid*128+tid];
__syncthreads();
for(int i=64;i>0;i/=2)
{
if(tid<i)
s_data[tid]=s_data[tid]+s_data[tid+1];
__syncthread();
}
if(tid==0)
B[bid]=s_data[0];
}
int main()
{
VecAdd<<<32,128>>>（A,B);
} ??????????
------------------------deviceQuery.cu---------------------------------------
#include<stdlib.h>
#include<string.h>
#include<stdio.h>
#include<cutil.h>
int main(int argc,char** argv)
{
int deviceCount;
CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));
if(deviceCount==0)
printf("there is no device supporting CUDA\n");
int dev;
for(dev=0;dev<deviceCount;++dev)
{
cudaDeviceProp deviceProp;
CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp,dev));
if(dev==0)
{
if(deviceProp.major==9999&&deviceProp.minor==9999)
printf("there is no device supporting CUDA.\n");
else if(deviceCount==1)
printf("There is 1 device supporting CUDA\n");
else
printf(" ");
}
//输出的cuda的设备数量。
CUDA C 语言
1.引入了函数类型限定符。用来规定函数实在host还是在device上执行，以及这个函数实在host还是在device
上调用。__device__ ,__host__,__global__.
2.引入了变量类型限定符。用来规定变量被存储在哪一类的存储器上。
3.引入了内置矢量类型。char4,ushort3,double2,dim3等，由基本的类型构成。
4.引入了4个内建变量。blockIdx和theradIDx用于索引线程块和线程，gridDim和blockDIm用于描述
线程网格和线程块的维度。warpSize用于查询warp中线程的数量。
5.引入了<<<>>>运算符。用于指定线程网格和线程块的维度，传递执行参数。
6.引入了一些函数：memory fence函数，同步函数，数学函数，纹理函数，测试函数，原子函数