Cuda learn record three

1. 页锁定 主机内存

通过malloc分配的主机内存是可分页的（pagable），而cudaHostAlloc要分配页锁定的主机内存，即固定内存。GPU可以直接内存访问。

但是，使用页锁定内存时，操作系统的虚拟内存的功能将失效，程序可能会快速的耗尽系统内存，这一方面可能导致程序因内存不足而运行失败，另一方面，也会使其他正在使用的程序的效率受损。

固定内存是一把双刃剑。使用页锁定时，不需要该内存时，及时的释放掉该内存。cudaFreeHost（）

当系统的PCIE总线带宽受到限制的时候，使用页锁定内存可以有效的加速程序。

2. CUDA 流

以异步方式只能对页锁定内存进行复制操作。

异步操作只能保证会在下一个被放入流中的程序之前执行，但不能保证返回时已经运行结束。

 3. cuda中很少使用2维以上的指针，这与cuda本身在内存中的复制有关。

但是我们可以使用一维线性的内存空间来模拟二维甚至三维的数组。

下面这个地址详细的介绍了二维和三维类型的数组的申请方法。

http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015554.html

或者是，

http://blog.csdn.net/augusdi/article/details/12187091

4. 对于C语言中memset（）函数的认识。

该函数的原型是void *memset(void *s, int ch, size_t n);

但是要注意到的一点是，该函数对于非字符型的数组，只能赋0或者-1，其他的赋值都会出问题。这是因为该函数采用的赋值的方式是逐个字节的赋值。

5. 考察了C语言用指针生成二维动态数组的方法。

template<typename T>
T** createArray(T b, int m, int n)
{
    T ** p = (T **)malloc(m * sizeof(T*));
    assert(p != NULL);
    T* q = (T *)malloc(m*n * sizeof(T));
    assert(q != NULL);
    memset(q, (T)0, m*n * sizeof(T));
    //printf("%d@@", sizeof(q));
    for (int i = 0; i != m; ++i) {
        p[i] = i*n+q;
    }
    for (int i = 0; i != m; ++i) {
        for (int j = 0; j != n; ++j) {
            p[i][j] = 2;
        }
    }
    return p;
}

template<typename T>
void deleteArray(T** p) {
    free((void*)p[0]);
    p[0] = NULL;
    free((void*) p);
    p = NULL;
}

//下面是另一种实现的方法

template<typename T>
void allocMatrix(T*** pointer, int m, int n)
{
(*pointer) = (T **)malloc((m + 1) * sizeof(T *));
if ((*pointer) == nullptr) {
printf("内存空间不足，无法申请二维数组！ ");
exit(1);
}
for (int i = 0; i <= m; ++i) {
(*pointer)[i] = (T*)malloc((n + 1) * sizeof(T));
if ((*pointer)[i] == nullptr) {
printf("内存空间不足，无法申请二维数组！ ");
exit(1);
}
}
}

template<typename T>
void deleteMatrix(T ** &p, int m) {
for (int i = m; i >= 0; --i) {
free(p[i]);
p[i] = nullptr;
}
free(p);
p = nullptr;
}

实现了一段在host端使用二维数组，在device端使用一维数组的程序

    int temp = 0;
    //int *host_a =(int *) malloc(5 * 5 * sizeof(int));
    int **host_a = createArray(temp, 5, 5);
    for (int i = 0; i != 5; ++i) {
        for (int j = 0; j != 5; ++j) {
            host_a[i][j] = 1;
        }
    }
    int *dev_a = NULL;
    int *h_a = host_a[0];
    for (int i = 0; i != 5; ++i) {
        for (int j = 0; j != 5; ++j) {
            cout << h_a[i * 5 + j]<<" ";
        }
        cout << endl;
    }
    HANDLE_ERROR(cudaMalloc((void**)& dev_a, 5 * 5 * sizeof(int)));
    HANDLE_ERROR(cudaMemcpy( dev_a, *host_a, 5 * 5 * sizeof(int), cudaMemcpyHostToDevice));

    dim3 threads(5, 5);
    kernel << <1, threads >> > (dev_a);

    HANDLE_ERROR(cudaMemcpy(*host_a, dev_a, 5 * 5 * sizeof(int), cudaMemcpyDeviceToHost));

    for (int i = 0; i != 5; ++i) {
        for (int j = 0; j != 5; ++j) {
            cout << h_a[i * 5 + j] << " ";
        }
        cout << endl;
    }

    free(host_a);
    HANDLE_ERROR(cudaFree(dev_a));