最近终于等到一个比较好的时机可以摸一摸心心念念的CUDA了，不过手边唯一的一个NV显卡是游戏本上的1050Ti。为了保持 娱乐的纯洁性，游戏本上也只有Windows系统。课题组里的GPU服务器资源有点不好意思要，毕竟现在只是在玩一玩的阶段。最后只有考虑弄个外面的GPU云服务器了。

炼丹七天乐.jpg
炼丹七天乐.jpg

  春节的时候，AutoDL弄了个“炼丹七天乐”，个人感觉体验不错，他们的服务器是按分钟计费的，虽说价格还算实惠，但是再实惠也没有白嫖爽呀，于是我大踏步地投入了Colab的怀抱。Colab本来是跑Jupyter Notebook的，方便炼丹师摆弄他们的Python模型，但是实际上就是一个container可以挂载GPU设备，于是使用一些魔法符号也能直接操作shell。但是在 hello world 阶段出现了一些问题，在此记录一下。

  Colab分配的是K80设备，计算能力为3.7，毕竟是免费嘛。但是默认使用cuda11，使用nvcc编译的时候，如果不加-arch的flag，指定对象设备的加工，会默认采用sm_52，也就是计算能力为5.2的架构。这就导致编写的kernel无法顺利在GPU上执行，而CUDA10及以下的版本，都是默认采用兼容度最大的架构作为编译目标架构。

  此外，CUDA并不主动报出Kernel执行的信息，也就看不到对应的错误。这里参考了stackoverflow的相关方法，通过定义一个macro去处理runtime信息。

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

  使用的话，就是在Kernel调用后立即拉取一下：

HelloKernel<<<1,4>>>();
gpuErrchk( cudaGetLastError() );
gpuErrchk( cudaDeviceSynchronize() );

  另外的话，Colab上面panel会时不时出来几条小狗，还挺可爱的，哈哈哈。

dog_in_colab.png
dog_in_colab.png