cuda error与性能提升笔记

记录一下出现的错误，防止之后出现了还得一遍遍排查。

一、 CUDA error

1. Error 700 非法内存访问

排查后，发现kernel函数里if语句写错了，少了等号，应该是

if (n>=N) return;

2. 内存访问错误

* 可能是因为分配内存的时候错了

// 错误的情况，这样的分配只分配了num个8字节（地址的大小）
SAMPLEPOINT* pSample;
CHECK(cudaMallocManaged(&pSample, sizeof(pSample)*num));

// 修正
SAMPLEPOINT* pSample;
CHECK(cudaMallocManaged(&pSample, sizeof(SAMPLEPOINT)*num));

* 先分配内存，再引用指针

// 不可以！！
PHASESPACE* phsp = rtplan.phsp;
CHECK(cudaMalloc(&rtplan.phsp, sizeof(PHASESPACE)*N));

// 正确
CHECK(cudaMalloc(&rtplan.phsp, sizeof(PHASESPACE)*N));
PHASESPACE* phsp = rtplan.phsp;

3. 串行的时候没问题，并行的时候结果出错

串行的时候没问题（一般来说都是检查第一个结果），并行的时候结果出错。

原因：不同线程之间内存有交叠（内存分配错误或者数据溢出），或者浮点误差。

* 可能是因为指针用错了，一定注意使用的是该线程下应该使用的指针。

* 可能是内存分配错误，注意sizeof(type)是不是对了。

* 串行和并行在计算时结果不一样，串行是0.0001，并行很可能就-inf了。万一是算一个索引，索引由于浮点误差，成负数、inf了，就会报错非法内存访问。

// Input: DATASET* dataset;
size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
DATASET* dataset_ptr = dataset + tid;

// note!
float a = dataset_ptr->a;
// error
float a = dataset->a;

4. F5进行调试时，运行了一个核函数发现变量栏的内存访问都变红且'???'。看看是不是忘记同步了！

test_kernel<<<1,1>>>(a);
CHECKKERNEL;
CHECK(cudaDeviceSynchronize()); // 长时间不写就总忘记这忘记那的

5. 有关bool数组的结果出错

可能不是因为程序错了，是没有对bool数组进行初始化（全部置0或1）

6. 无法解析的外部函数

之前编译的好好的，忽然出现这个错误，可能是脑子抽了修改了函数的定义但没有修改对应的声明！比如把void func(const string& filename), 改成void func(string &filename)，编写的时候vs根本不会报错，但编译不会通过。

二、性能提升

按照性能提升最大的顺序排列，不一定适合所有代码。

1. 数组解引

当N很大时真的可以快很多，可能是因为我的代码里有很多关于数组的计算。

for (int i = 0; i < N; i++) {
    A[i] = i;
}


// modified
for (int i = 0; i < N; i++) {
    *(A + i) = i;
}

2. 循环展开

for (int i = 0; i < 3; i++) {
    A[i] = i;
}

// modified
A[0] = 0;
A[1] = 1;
A[2] = 2;

to be continued..