在GPU中 设备函数的参数可以是指向某存储区域的指针吧？也可以带返回值吧？譬如返回数组最大值的位置：
device int maxpos(float c, int m)
{
int maxid=0, k;
float val, maxval = *c;
for (k=1; k<m; ++k) {
val = c[k];
if (val > maxval) {
maxval = val;
maxid = k;
}
}
return maxid;
}
程序这样写应该没有问题吧，求指教。
当然，设备函数和主机函数没什么区别，只是要求传进来的参数在显存上。
LZ您好:
1.device 函数可以使用指针参数，可以有返回值。指针需指向GPU端的存储空间。
您这样写语法上无问题，可以用。
但是__device__ 函数和kernel一样是每个线程都跑一份的，所以您还需要注意这点，让多个线程分担您的任务，以及注意合并访问等。
大致如此，供您参考。
祝您好运！
针对第三点 让多个线程分担任务和合并访问具体是什么意思？ 有相应的例程吗？
LZ您好
这个要求其实和对kernel的要求是一样的。
例如您有n个线程，要写入8n个数据，那么每个线程写8个数据而不是一个线程写入8n个数据，这个是让多个线程分担任务。
以及在写入的时候，需要考虑一个warp的行为，此时一个warp内部的线程写入一段连续的空间内，大致上是合并访问的。
大致如此，祝您好运~
版主你好，假设我的global函数要调用很多device函数，是不是应该在global的.cuh文件中包含这些device函数？
LZ您好：
1：如果您的__device__函数和__global__函数在同一文件中，那么直接用即可。
2：如果不在同一文件中，那么请参考如下链接中横扫斑竹给出的建议：
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6344
祝您好运~
那对于__device__函数调用__device__函数，也是同样的做法吧？
LZ您好，是一样的。
祝您编码顺利~
版主您好，
看到global函数里调用device函数时没有用<<<gridSize,blockSize>>>指定并行方式，那是不是global函数里每个线程都会把这个device函数执行一遍？在global函数调用device时是否可以指定<<<gridSize,blockSize>>>，那device函数里面也有threadIdx等内置变量吗？
谢谢斑竹大人！
您好：
1：__device__函数是每个线程都跑一份的。
2：如果您使用计算能力为SM3.5的kepler核心的GPU，那么可以在kernel里面调用kernel，此时作为调用kernel是可以指定<<<>>>形状的。但是调用__device__函数，只能是一个线程跑，并无法设定<<<>>>形状。
3：__device__函数内依然可以使用threadIdx.x这样的内建变量来确定线程编号，以及，如果您在调用__device__函数之前已经利用内建变量生成了您的线程编号，那么也可以直接作为__device__函数的参数传进来使用。
大致如此，祝您好运~
明白了，多谢斑竹:loveliness:
不客气的，欢迎您常来论坛，祝您好运~