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.之后会发生什么？我的意思是现在有两个并发的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]++;
}

执行并行发生,就像普通代码一样.如果您将内存竞争设计为函数,则没有可以节省您的序列化机制.