CUDA编程debug与release模式结果不一致的

最新推荐文章于 2024-04-28 11:19:39 发布

去看樱花。

最新推荐文章于 2024-04-28 11:19:39 发布

阅读量274

点赞数

分类专栏： CUDA 文章标签： c++

本文链接：https://blog.csdn.net/weixin_42516268/article/details/132102665

版权

CUDA 专栏收录该内容

1 篇文章 0 订阅

订阅专栏

咱就是说咱找这个问题找的花儿都谢了。先说结论：

CUDA的默认设置对浮点运算非常不友好，精度很低，而且release下cuda为了速度快竟然牺牲精度！！！！。（可以解决）

修改方法：

项目属性 --》配置属性--》cuda c/c++ --》Command Line 界面中的其他选项中添加：-fmad=false。fmad(floating-point multipies and add/subtracts into floating-point multiply-add operations),大概就是会把乘法跟加法融合到一起做一些省略型的计算的意思。禁用FMAD整个编译单元的优化，有损性能，但是能提高精度。

记录一下过程：

关于多grid与多block间线程同步问题，核函数如下：

__global__ 
	void t_approxgrid(TDOSE *pDose，...XXXXXXX){

pDose[z * size.y * size.x + y * size.x + x] += dose;

}

入口函数主要关于同步方面内容如下：

	// 在GPU上分配内存，并将剂量数值拷贝过去
	size_t lens = size.x * size.y * size.z * sizeof(TDOSE);
	checkCudaErrors(cudaMalloc((void **)&output,lens));
	checkCudaErrors(cudaMemset(output,0, lens));	
	checkCudaErrors(cudaMemcpy(output,pDose,lens, cudaMemcpyHostToDevice));

	// 定义时间记录事件
	cudaEvent_t start, stop;
	checkCudaErrors(cudaEventCreate(&start));
	checkCudaErrors(cudaEventCreate(&stop));
	checkCudaErrors(cudaEventRecord(start, 0));

        t_approxgrid<<<gridSize, blockSize>>>(output,...XXXXXX);
 
	checkCudaErrors(cudaEventRecord(stop, 0));
	checkCudaErrors(cudaEventSynchronize(stop));
	checkCudaErrors(cudaEventSynchronize(start));
	//cudaDeviceSynchronize();
 
	float elaptime;
	cudaEventElapsedTime(&elaptime,start,stop);

	checkCudaErrors(cudaMemcpyAsync(pDose,output,lens, cudaMemcpyDeviceToHost));		
	checkCudaErrors(cudaFree(output));

问题：

感觉pDose[z * size.y * size.x + y * size.x + x] += dose应该有线程间争抢内存，内存存储错误的问题，但是测试了一些并没有影响。保险起见改成了atomicAdd(&pDose[z * size.y * size.x + y * size.x + x],dose);

结果一致（哭），但是的话呢，用同样的坐标z * size.y * size.x + y * size.x + x打印了一个累加计数器保存到本地，atomicAdd跟线程数量一致，+=操作比线程数量少（晕）