如何处理GPU上Error Number:700 an illegal memory access was encounter

【现象描述】

GPU上网络运行过程中出现Error Number:700 an illegal memory access was encounter

 

【原因分析】

出现该现象，在框架稳定的背景下基本上可以确定是网络中有算子踩显存，因此CUDA上报非法内存访问，错误码为700，可能原因如下：

1.算子计算过程中使用的size比申请的显存大了，导致访问越界。

2.由于GPU的算子执行是host下发到device上异步执行的，host使用了CUDA一些同步接口导致不是device的期望值出现非法内存。

 

【解决方法】

步骤1：由于GPU的算子执行是host下发到device上异步执行的，因此执行报错的地方不一定是真凶，大概率是前面的算子有问题，但是device是异步执行的，所以执行到后面才会报错。可以设置环境变量export CUDA_LAUNCH_BLOCKING=1表示阻塞式执行，也就是一个算子在device执行完成后，host才会下发一个算子到device上执行，这种完全同步执行方式下如果还是报700非法内存错误的话，配合算子执行日志就可以确定是当前执行算子的问题了。

案例代码：https://gitee.com/mindspore/mindspore/pulls/962

 

 

步骤2：执行步骤1的操作后问题不复现，也就是同步执行的这种方式下没有问题，基本上可以确定是有算子里依赖同步执行的结果，因为正常device算子执行是异步执行，所以拿的结果不是预期值，同步执行就掩盖了这个问题。出现这种情况，除了走读代码去确认是哪个算子里有使用同步操作接口外，一般可以通过二分法去加同步流来去定位（如果一个网络中有100个算子，可以在第50个算子加同步流，如果加了同步流没有复现，说明就前面的某个算子有问题了），同步流添加方法：在GPUDeviceContext::LaunchKernel函数中调用DoLaunchKernel后面增加SyncStream()。

案例代码：

```

void FastTopK(const int outer_size, const int inner_size, const T *input, const S *k, T *output, S *output_index,
              const T init_K, cudaStream_t stream) {
  int block_num_limit = outer_size < 128 ? outer_size : 128;
  S k_cut = 0;
  cudaMemcpy(&k_cut, k, sizeof(S), cudaMemcpyDeviceToHost);
  if (k_cut > inner_size) k_cut = inner_size;

```

---》此处使用了cudaMemcpy期望从device同步到host，是个同步接口，但是实际上此次device的还未执行完，因此不是期望的值，导致后面执行逻辑都是错乱的，出现非法内存访问。

 

【建议与总结】

1.算子实现逻辑中针对size需要check，确保访问不越界。

2.由于GPU的算子执行是host下发到device上异步执行的，尽量避免使用CUDA同步接口，如果使用的话，需要确保数据是正确的。