int **d_ws;
[indent]cudaMalloc((void**)(&d_ws), MAXCOUNTMAXCOUNTsizeof(int));[/indent][indent]cudaMemcpy((void*)(d_ws), (void*)(h_ws), MAXCOUNTMAXCOUNTsizeof(int*), cudaMemcpyHostToDevice);[/indent]
如上面所示，可是失败了。
现在请问有标准的拷贝二维函数的函数吗？
目前有个想法是把 h_ws二维函数转化成1维，拷贝到device上面，再转化成二维来执行操作。
LZ您好：
您确定您是要拷贝“二维函数”？
如果您是要拷贝二维数组，那么可以将其一维化，然后安排好寻址等操作即可。
您也可以参考如下帖子：
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=7056&extra=page%3D1
祝您好运~
我继续补充下ICE大在2#的建议：
前提：楼主您的确需要的是2维数组，而不是“2维函数”。
在这个前提达成共识下：
（1）楼主您的定义有错误，
您应当定义成ICE建议的int *d_ws(一维化后的），
或者定义成：int (*d_ws)[MAXCOUNT];
但绝不是您的int **d_ws;(这实际上是个一维指针，指向的是一些一维指针构成的缓冲区）。
修改后，您可以继续使用cudaMalloc分配您的MAXCOUNT * MAXCOUNT个元素的缓冲区。
在您修正了定义问题后，您还有个错误：
(2)您cudaMemcpy()的时候，使用了sizeof(int *), 根据您的上下文，您需要的是sizeof(int)
请您考虑本文和ICE上文的建议。修改代码。
感谢来访。
以及，需要补充的是：
如果楼主的MAXCOUNT * MAXCOUNT个元素的数组，
在MAXCOUNT * sizeof(int)无法被128整除的时候（在1.x和2.x上），或者无法被64整除的时候（在3.x上），建议您不要使用cudaMalloc进行一维化创建。
此时推荐cudaMallocPitch进行分配，此时将有可能会提升部分性能。
（请注意：cudaMallocPitch虽然可能提升性能，但也可能会增加您的显存使用量，请考虑此2因素，酌情决定是否需要使用cudaMallocPitch)
感谢您的来访。
数组大小都是我来控制，尽量让它在128倍数就行啦。那我就用一维的来表示好了。谢谢版主及ICE了。但是还有一个问题： 我目前有256个点，我想让这256个点分两次执行，也就是kernel<<<1,128>>>(…)这样声明的，我目前想让前k个点（第0个，第1个，…第k-1个）执行A操作，后k个点（第k个，第k+1个，…第255个）执行B操作，中间的执行C操作。怎么能区分第一次执行128里面的前k个点，以及第二次执行的后k个点。我的k很小，一般也就是4-6而已。如果用tid来表示这个点，怎么区分第一次执行的tid=1的点，跟第二次执行tid=1得点。
首先说，楼主您没认真看我的回复，我说的是*sizeof(int)后不是128或者64的倍数，而不是按元素数计算的宽度。
其次，没看明白楼主您的第二个问题是什么意思？
如果您需要启动2次kernel，同时能够区别这2次kernel, 您可以直接写2个kernel的。
或者如果您不想写，也可以单独设置一个参数，例如第一次用：
<<<>>>(1, …); 启动。
<<<>>>(2, …); 启动
然后kernel函数内部进行判断。
这样也可以。
请您判断这是否是您需要的解决方案，如果不是，请您重新用清晰的语言描述您的问题是什么，然后论坛将再次给出解答。
感谢来访。
噢，通过block来控制区分每次的tid就行啦。谢谢版主。另外上面的那个没看清楚，只要保证对齐就可以了，对吧！
不对齐也没关系的，依然可以正确的运行的（你懂的）。
恩，好的，谢谢版主
您客气了，感谢您的来访，欢迎您经常来提问。