CUDA C语言对C语言的扩展主要包含以下几个方面。
1.函数限定符
__device__：声明在设备上执行的函数。该函数无法被CPU端调用，只能由GPU端程序调用，即只能被__device__或__global__声明的函数调用。
__global__：声明的函数称为kernel函数。该函数只能被CPU端调用，执行在GPU上。Kernel函数类型必须是void,即返回类型必须为空。
__host__：声明在主机执行的函数，仅可在CPU端调用。一般情况可省略，只有该函数同时存在被设备端和主机端同时调用的情况需要添加该限定符，且无法与__global__联用。
2.变量限定符
__device__:声明在设备上的变量，该变量位于global memory(参见3.3节GPU存储体系)，只能在设备端使用，是全局变量，无须也不能在函数参数表中出现。
__constant__:声明在常量存储中的变量，只能在设备端使用。此变量一般情况下是只读的，只能通过特定方式进行修改（详见16.6节）。该变量是全局变量，无须也不能在函数参数表中出现。
__shared__:声明在共享存储中的变量，仅供block内所有thread共享访存，退出kernel函数后失效。该变量无法初始化，一般声明在kernel函数中。
3.内置数组变量
char2、uchar2、char3、uchar3、char4、uchar4、short2、ushorta2、short3、ushort.3、short4、ushort4、int2、uint2、int3、uint3、int4、uint4、long2、ulong?2、long3、ulong3、long4、ulong4、f1oat2、float3、loat4、double2:结构体数据.
4.内建变量
gridDim:指定grid维度，类型为dim3。
blockDim:指定block维度，类型为dim3。
blockIdx:指定grid内block索引号，类型为uint3。
threadIdx:指定block内thread索引号，类型为uint3。
warpsize:指定warp内thread数量，类型为int。
5.kernel调用
kernel函数调用就是调用声明为__global__的函数，必须在主机端调用。调用时需要指定线程维度，指定方式为<<Bs,Ts,Ss,Si≥>，其中Bs指定gid内block维度，类型为dim3;Ts指定block内thread维度，类型为dim3;Ss指定共享存储空间大小，类型为size_t;Si指定流索引号；其中Ss和Si可选。
下面给出了一个实例（详细代码见第6章）进行说明，其中，blocknum为grid内的block数量，threadnum为block内的thread数量。vector_add_gpu_3<<<blocknum, threadnum>>>(d_a,d_b,d_c,n);
6.特殊函数和内建函数
CUDA C中引入了大量专用的特殊函数和内建函数以实现特定功能，比如同步函数、数学函数、纹理函数、测时函数、原子函数、存储栅栏函数等。
