CUDA 踩坑

最近踩的坑

1.  声明using namespace nvcuda的时候，编译要加上-arch=sm_70选项。

2. 给half类型变量赋常量的时候必须加小数点，不然赋值完就是0。

3. kernel的时候用到动态shared memory，结果在给shared memory赋值的时候一直报错

“an illegal memory access was encountered”
原因：使用动态shared memory时必须在调用核函数的时候声明shared memory的大小，并且大小要保证kernel不会访问声明大小以外的地址。

4. WMMA操作不知道为什么直接使用fragment.x[]访问单个元素的值并不对（本来想通过这样将一个fragment中的元素复制到另一个fragment，最后还是采取先store到shared memory再load的方法）。

5. 要使用 cooperative_groups的时候编译选项需加上 -rdc=true。

6. 访问global memory的时候如果超过了malloc的分配空间居然编译和运行时都不会报错，所以一定要检查地址和指针的运算。

7. 如果一次计算的结果不会被存入shared，global memory， 那计算就会被编译器优化掉, ggwp。对于一次tenor core operation如果计算结果中有一个数据有存入memory的操作那这次tensor core operation就不会被优化。