CUDA:从内核调用 __device__ 函数

Question

我有一个内核，它在 if 语句中调用 device 函数。代码如下:

__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。每个函数调用都是串行执行的吗？那么它们的行为就像两个内核函数调用？

Answer 1

CUDA 实际上默认内联所有函数(尽管 Fermi 和较新的架构也支持具有函数指针和实际函数调用的适当 ABI)。所以你的示例代码被编译成这样的

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

执行是并行发生的，就像普通代码一样。如果您将内存竞争设计到函数中，则没有序列化机制可以拯救您。