CUDA：从内核调用__device__函数

Question

我有一个内核调用if语句中的设备函数函数。代码如下：

__device__ void SetValues(int *ptr,int id) 
{ 
    if(ptr[threadIdx.x]==id) //question related to here 
      ptr[threadIdx.x]++; 
} 

__global__ void Kernel(int *ptr) 
{ 
    if(threadIdx.x<2) 
     SetValues(ptr,threadIdx.x); 
} 

在内核线程0-1中同时调用SetValues。那之后会发生什么？我的意思是现在有2个并发呼叫SetValues。每个函数调用是否都是串行执行的？所以他们的行为就像2个内核函数调用一样？

Answer 1

CUDA默认内联了所有函数（虽然Fermi和更新的架构也支持带有函数指针和实函数调用的正确ABI）。所以，你的示例代码被编译到这样的事情

__global__ void Kernel(int *ptr) 
{ 
    if(threadIdx.x<2) 
     if(ptr[threadIdx.x]==threadIdx.x) 
      ptr[threadIdx.x]++; 
} 

执行并行发生，就像正常的代码。如果你将一个内存竞争设计成一个函数，那么没有序列化机制可以为你节省。