cudaMallocPitch和cudaMemcpy2D的设备端内存分配

cudaMallocPitch和cudaMemcpy2D的设备端内存分配

cuda调试中遇到了一个bug，调bug的过程中发现对cudaMallocPitch的内存分配理解错了，记录如下：

---

bug场景

.cpp中定义了几个类成员变量：

 float (*f_vertices)[3] = nullptr;
 float (*f_normals)[3] = nullptr;
 float (*dv_vertices)[3] = nullptr;
 float (*dv_normals)[3] = nullptr;

其中，f开头的两个成员变量作用是存储在内存中开的一段数组空间，dv开头的两个成员变量是存储在显存中开的一段数组空间。

开显存方式如下：

m_gpuFunc.send_to_gpu<float(*)[3], float*, float, 3>(model.f_vertices, model.dv_vertices, model.Vsize);
m_gpuFunc.send_to_gpu<float(*)[3], float*, float, 3>(model.f_normals, model.dv_normals, model.Vsize);
//上面是.cpp里的函数，调用下面.cu的函数
template<typename Datatypes1, typename Datatypes2, typename Eletypes, int num>
__host__ void gpuFunc::send_to_gpu(Datatypes1& h_data, Datatypes2& d_data, int height) {
    if(d_data != NULL)
        cudaFree(d_data);
    size_t pitch;
    cudaMallocPitch((void **)&d_data, &pitch, sizeof(Eletypes) * num, height);
    m_pitch = pitch;
    HANDLE_ERROR(cudaMemcpy2D(d_data, pitch, h_data, sizeof(Eletypes) * num, sizeof(Eletypes) * num, height, cudaMemcpyHostToDevice));
}

用这种方式开显存后，编译器虽然没有报错，但在gpu核函数中对dv_等变量用二维数组的[i][j]方式寻址总是不能正确得到结果。
贴个核函数的声明

__global__ void g_calcFaceDist(int (*&triangles)[3], float (*&points)[3],  size_t& triHeight, size_t& poHeight, float (*&facesNorm)[3], float *&facesDist);
//在该函数中调用triangles[i][j]得不到想要的值。

---

原因分析：

引用传值问题

首先是发现核函数不能引用传值，这个坑蛮大的，我也是调的生无可恋后偶然把引用删了，线程突然就不出错了，个人理解是__host__函数用的变量是存储在内存中的，__global__函数用的变量是存储在显存中的，显存中的引用（引用是个指针）指向的是显存中的空间（无效）而不是原来的内存中的空间，故寻址失败。

空间分配与组织问题

虽然线程不出错，但是对triangles这个数组采用[i][j]寻址总是得到0这个结果，打印了所有输出，发现大部分都是0，于是想到可能是寻址方式不对。

仔细看了下cudaMallocPitch和cudaMemcpy2D的文档，发现这个malloc方式是开多段内存对齐的行，行间偏移量（行大小）为pitch（我的这段程序中是512），两个函数声明如下：

cudaError_t cudaMallocPitch(void** devPtr, size_t* pitch, size_t widthInBytes, size_t height)；
cudaError_t cudaMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);

cuda根据所传的widthInbytes（行大小）来生成一个更大的数值为512倍数的pitch，在显存中开辟行偏移量为pitch的数组。在核函数中访问这个数据对象时，如想要访问第i行第j列，则索引为
[i * pitch/sizeof(datatypes) + j]
其中datatypes为数组中数据的类型。

由以上特性，问题就显而易见了，我的核函数中使用的数据是一个float(*)[3]类型的，即编译器将其看成按3分组的指针数组，但是在显存中他存的其实是按128分组的数据，所以float(*)[3]类型的triangles与他存的数据根本那就不能对齐，128甚至都不是3的倍数，所以存在显存中的数据还是使用单独的指针类型比较好，访问灵活。

其实从cudaMalloc2D的声明也能看出来，第一个参数就是一个void**的地址，我原先写的float(*)[3]取完地址以后都是void***了，说明这个函数本身就建议传入一个单独的指针。

遂把这些数据类型及函数模板全都改了：

//变量
float (*f_vertices)[3] = nullptr;
float (*f_normals)[3] = nullptr;
float *dv_vertices = nullptr;
float *dv_normals = nullptr;

//调用
m_gpuFunc.send_to_gpu<float(*)[3], float*, float, 3>(model.f_vertices, model.dv_vertices, model.Vsize);
m_gpuFunc.send_to_gpu<float(*)[3], float*, float, 3>(model.f_normals, model.dv_normals, model.Vsize);

//模板描述
template<typename Datatypes1, typename Datatypes2, typename Eletypes, int num>
__host__ void gpuFunc::send_to_gpu(Datatypes1& h_data, Datatypes2& d_data, int height) {
    if(d_data != NULL)
        cudaFree(d_data);
    size_t pitch;
    cudaMallocPitch((void **)&d_data, &pitch, sizeof(Eletypes) * num, height);
    m_pitch = pitch;
    HANDLE_ERROR(cudaMemcpy2D(d_data, pitch, h_data, sizeof(Eletypes) * num, sizeof(Eletypes) * num, height, cudaMemcpyHostToDevice));
}

//核函数声明
__global__ void g_calcFaceDist(int *triangles, float *points,  size_t triHeight, size_t poHeight, float *facesNorm, float *facesDist)；

这样写就一切正常了，

还有个事就是函数模板的声明与实现如果放在两个文件中，需要在实现文件里显式的对用到的类型进行实例化，否则编译会报错，原因是函数模板会经过两次编译，第一次编译时在实现文件里没有找到示例化的模板函数的话就会认为这个函数没有被实现。
显示实例化例子如下：

template void gpuFunc::send_to_gpu<float(*)[3], float*, float, 3>(float(*&)[3] , float*&, int);
template void gpuFunc::get_from_gpu<int(*)[3], int*, int, 3>(int(*&)[3], int*&, int);