__device__ function参数传递讨论

手册上说__global__function parameters are passed to the device:via shared memory and are limited to 256 bytes on devices of compute capability 1.x; via constant memory and are limited to 4kb on devies of compute capability 2.0
那么__device__ function 的参数是如何传递的呢?对于tesla架构来说,它默认是内联的,可以认为它只是个形式上的函数,而实际上并没有像真正的函数那样传递参数;但是 feimi架构支持非内联的__device__function,那么此时函数的参数是如何传递的?也是通过shared memory或者是constant momory吗??
大家讨论一下吧,有哪位大牛能告诉我吗??谢谢!