CUDA采坑之路

1.cudaMalloc来回拷贝东西的时候一定要注意内存大小的问题，不然的话可能会出问题。

2.CUDA主机函数中可以使用c的语法，可以使用printf进行调试。

3.printf有的时候可以起到__syncthreads()的作用，所以有的时候如果加了printf能用的话，很可能是忘了加同步了。

4.在改变任何所有的核函数统一用到的变量之前一定要加__syncthreads()，如下面的例子中，因为每次i更新后所执行的操作都必须要求i为上一次的值得时候的操作都执行完成，如果不加同步，就会造成执行顺序无法控制，最终造成结果的错误。

 for (int i = blockDim.x / 2; i >= 1; i /= 2) {
        if ( threadIdx.x < i) {
            c_cache[ threadIdx.x] += c_cache[ threadIdx.x + i];
            __syncthreads();
        }
    }

5.double有的时候比较占空间，所以，炸了的话，换float就行了。

6.对于texture一定要在核函数（使用到它）之前定义。同时对于texture的理解，把a（texture变量）和b（全局变量）绑定了候，每次在执行核函数的时候会把b拷贝到a中，在这次核函数执行的过程中，对于b中的内容仍然可以修改，但是a中的内容是不变的。然后对于下次调用核函数的时候，a会再次把b拷贝，这次的就是上次对b修改后的内容。示例代码

#include "cuda_runtime.h"
#include"cuda.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include <iostream>
using namespace std;
#define N 10
texture<float> test;
#define mins(a,b)(a<b?a:b)
#define threadsperblock 16
const int blocknum=mins(int((N +threadsperblock)/ threadsperblock), 32);

__global__ void kernel(float *a, float *b)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    a[id] = a[id] + b[id];
    b[id] = tex1Dfetch(test, id);
    printf("now a[%d] %f  texture[%d] %f\n", id, a[id], id, tex1Dfetch(test, id));   
}

int main()
{
    float a[N], b[N];
    float* da, * db;
    for (int i = 0; i < N; i++) {
        a[i] = i;
        b[i] = i *2;
    }
    cout << cudaMalloc((void**)&da, N * sizeof(float)) << endl;;
    cout << cudaMalloc((void**)&db, N * sizeof(float)) << endl;;
    for (int z = 0; z < 2; z++) {
        cout << cudaMemcpy(da, a, N * sizeof(float), cudaMemcpyHostToDevice) << endl;
        cout << cudaMemcpy(db, b, N * sizeof(float), cudaMemcpyHostToDevice) << endl;

        cudaBindTexture(NULL, test, da, N * sizeof(float));
        kernel << <blocknum, threadsperblock >> > (da, db);

        cudaMemcpy(a, da, N * sizeof(float), cudaMemcpyDeviceToHost);
        cudaMemcpy(b, db, N * sizeof(float), cudaMemcpyDeviceToHost);

        for (int i = 0; i < N; i++) {
            cout << a[i] << "," << b[i] << endl;
        }
        cout << endl;
    }
    cudaFree(da);
    cudaFree(db);
    cudaUnbindTexture(test);
}