3.内存管理：
Device上的内存可以被分配成线性的，也可以分配为CUDA的数组形式的。CUDA的内存可以为1维，2维，还有3维(2.0版本)。内存的类型有unsigned8，16或者32位的int，16位(只有driver API可以做到)float，32位的float。这里分配的内存也只能通过kernel里面的函数通过处理纹理的方法来处理。这个地方也是GPU的历史原因了，以前都是处理图像的，所以这里叫纹理。……叫啥都是别人取得名字 - -!在计算机里面不就是内存嘛 - -!
Host的runtime的运行库也提供按照page-locked的内存管理的函数，page-locked的内存要比pageable方式快很多。好的东西往往比较少～page-locked就是很稀少的。如果通过减少分配pageable的内存来分配多的page-locked内存，系统需要的分页内存就少了，这也就会让系统的性能降低了。所以在处理这块的时候要合理。
Runtime API：
使用 cudaMalloc() 或者 cudaMallocPitch() 来分配线性内存，通过cudaFree()释放内存.
下面是分配一个大小为256 float数组的方法:
float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(float));
在使用2D数组的时候最好用cudaMallocPitch()来分配，在guide的第五章在讲到内存之间的调度的时候，就会看到他的好处。下面是一个分配一个大小为width×height 2D float数组的例子:
float* devPtr;
int pitch;
cudaMallocPitch((void**)&devPtr, &pitch,
width * sizeof(float), height);
myKernel<<<100, 512>>>(devPtr, pitch);
// device code
__global__ void myKernel(float* devPtr, int pitch)
{
for (int r = 0; r < height; ++r) {
float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c) {
float element = row[c];
}
}
}
CUDA 的数组方式，需要用 cudaMallocArray()和cudaFreeArray(). cudaMallocArray()又需要cudaCreateChannelDesc()来管理，这个其实可以在第guide的第五章里面可以看到，我们后面也会详细的介绍内存的调度和管理，和传统的GPU的内存方式不一样的地方.
分配 width×height 32位float的CUDA array例子:
cudaCreateChannelDesc<float>();
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaGetSymbolAddress用来在全局中定位一个数组的位置，然后cudaGetSymbolSize()来回忆他分配的时候大小----如果有全局编程或者多线程编程经验的，或者用过几个函数同时处理一个数据的经验，都会为了把数据的独立性弄出来，不能让数据和函数耦合太大，一般都不会让函数直接牵扯上数据，只是在函数处理的时候重新定位数据----这地方有点绕～～
下面是一些例子，内存之间的拷贝，回想一下有几种内存～linear的有两个函数可以分配的，还有CUDA array的内存:
width * sizeof(float), height,
cudaMemcpyDeviceToDevice);
The following code sample copies some host memory array to device memory:
float data[256];
int size = sizeof(data);
float* devPtr;
cudaMalloc((void**)&devPtr, size);
cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);
从host上面拷贝内存到device的constant上面:
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
Driver API：
cuMemAllocPitch()被推荐来作为2D的数组分配函数，会在内存对齐方面做一些check～然后保证在处理数据(像cuMemcpy2D()这样的拷贝函数)的时候达到最优的速度。下面就是一个float的width×height的2D的数组，并在循环里面处理的例子: ----(这个地方扩展提两个东东：一个是内存对齐，做程序优化的时候，或者处理网络问题的时候，都会遇到这些问题，内存对齐是一个普遍存在的问题，像SSE这样的编程的时候也是要求内存对齐的，看以后要是有机会以单独讲解内存对齐的问题：)第二个是CUDA的内存访问的问题，就是很多朋友都会在写kernel的时候，搞不明白里面的threadid，block id和传进来的内存的关系，这个地方必须要搞清楚的;内存和线程是CUDA编程必须搞明白的两个概念，不然到时候就会很混乱。其实看看前面的章节，应该能看明白的，不明白就用手画图～要是感觉还是有点不太清楚，也不要担心，接下来的章节会单独把一些难懂的问题更详细的讲解。)
CUdeviceptr devPtr;
int pitch;
cuMemAllocPitch(&devPtr, &pitch,
width * sizeof(float), height, 4);
cuModuleGetFunction(&cuFunction, cuModule, “myKernel”);
cuFuncSetBlockShape(cuFunction, 512, 1, 1);
cuParamSeti(cuFunction, 0, devPtr);
cuParamSetSize(cuFunction, sizeof(devPtr));
cuLaunchGrid(cuFunction, 100, 1);
// device code
__global__ void myKernel(float* devPtr)
{
for (int r = 0; r < height; ++r) {
float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c) {
float element = row[c];
}
}
}
cuArrayCreate()和cuArrayDestroy()来创建和释放CUDA array类型的数组。下面是一个 width×height 32-bit float类型的CUDA array 分配的例子：
desc.Format = CU_AD_FORMAT_FLOAT;
desc.NumChannels = 1;
desc.Width = width;
desc.Height = height;
CUarray cuArray;
cuArrayCreate(&cuArray, &desc);
这个也是几个内存之间拷贝的例子:
memset(©Param, 0, sizeof(copyParam));
copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY;
copyParam.dstArray = cuArray;
copyParam.srcMemoryType = CU_MEMORYTYPE_DEVICE;
copyParam.srcDevice = devPtr;
copyParam.srcPitch = pitch;
copyParam.WidthInBytes = width * sizeof(float);
copyParam.Height = height;
cuMemcpy2D(©Param);
拷贝host上面的内存到device上面：
int size = sizeof(data);
CUdeviceptr devPtr;
cuMemAlloc(&devPtr, size);
cuMemcpyHtoD(devPtr, data, size);
