搞了一星期了，kernel内部总是调不对。
具体来讲，我做的kernel所实现的功能是用多重嵌套循环来寻找最优参数。首先从GMEM上读取数据到SMEM上，由于需要参数很多，所以把SMEM相应地分成了几部分。这里面既有动态数组（从kernel发射处给bytes值，然后kernel内部声明extern后做地址偏移），也有静态数组（直接在kernel内部声明），这个用着倒没觉得有啥问题。 出问题的地方是在循环内部读取动态SMEM地址。比如读取shared[ i+BLOCKSIZE]就实现不了，令j=i+BLOCKSIZE后读取shared[j]也不成功，只有读取shared[常量]才行。这个问题会不会是因为大量线程同时访问shared，产生了大量的bank conflict所导致？
另外，除了kernel<<<>>>里所需要的四个参数，传递给kernel的参数个数有限制么？
最后，我在kernel之后加了printf(“%s\n”,cudaGetErrorString(cudaGetLastError()))，多次执行kernel时第一次会显示是no error，之后全部都是unknown error，但传递给kernel的参数没有任何的改变。而且之前可以用的一个子程序，放到kernel之后使用，竟然还会出现内存分配失败或者是内存释放失败导致bus error或者segmentation fault，这会是啥问题呢？真是要哭了
问题有点多，还望大家见谅！
另外即使出现严重的bank conflict，也只会是变慢，不会出错。报错应该是其他问题。
手机党先答复如上，请发一下代码，一会上PC看。
chaogege辛苦了。
（1）无法理解楼主的无法读取shared[i + BLOCKSIZE]，指的是什么？？因为如果楼主的i和BLOCKSIZE无论是常数还是变量，显然对数组寻址都是可以的。建议chaogege用普通的白话简单说一下。
（2）传递给kernel的参数个数是否有限制未能从公开途径得知可否，但kernel的参数总大小是有限制的，在常见的2.x和3.x上，是4kB。
（3）在回答这个问题前，需要和楼主达成如下2个共识（也是常识）。
（3-A) 如果你在一个GPU上下文中进行操作，一旦有一步出错，那么就会步步出错，除非摧毁重来(例如cudaDeviceReset())。
（3-B) cudaErrorUnknown一般代表着kernel执行了对无效显存地址的访问， 或者进行了不满足特定访问大小的对齐要求而出现。（该错误可以在nsight下被捕获，提示为access violation, 略有不同)
在楼主阅读过(3-A)和(3-B)的基础上，楼主可以理解为何出现上述问题了：
对序列：
kernel<<<…>>>(…);
auto r0 = cudaGetLastError();
kernel<<<…>>>(…);
auto r1 = cudaGetLastError();
…
中，r0是立刻得到的，无法反映kernel的运行进展中的问题。所以此时你找不到错误（除非你的GPU足够快，你的kernel足够小且挂的足够快）.
同时，因为一旦一步出错，你会步步重复得到这个错误，所以你又得到了以后的哪怕正确的可用的kernel也会出错的现象。
（4）您客气了。这个是我们应该做的。
预祝楼主春节愉快！