记录一下出现的错误,防止之后出现了还得一遍遍排查。
一、 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));
3. 串行的时候没问题(一般来说都是检查第一个结果),并行的时候结果出错。(可能是因为指针用错了,一定注意使用的是该线程下应该使用的指针)
// 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)
二、性能提升
按照性能提升最大的顺序排列,不一定适合所有代码。
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..