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个内核函数调用一样吗？

- scatman

Answer 1

CUDA默认情况下实际上会将所有函数内联（虽然Fermi及以上架构也支持使用函数指针和真正的函数调用），因此您的示例代码会被编译成类似于以下内容：

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

代码执行是并行的，就像普通代码一样。如果你在函数中设计了一个内存竞争，那么没有任何串行化机制可以拯救你。

- talonmies

---

我有一张 Fermi 显卡，这意味着该函数不是非内联的吗？因此，在 SetValues 中的 threadIdx.x 是所有线程而不仅仅是线程 0 和 1 吗？ - scatman

在 Fermi 中，默认情况下仍会内联函数。我想我现在明白你实际上在问什么了 - 就是像 threadIdx 这样的内置每个线程变量在 device 函数内的作用域是什么。我接近正确了吗？ - talonmies

是的，这是我的问题。但由于__device__函数是内联的，那么我猜线程的范围与使用<<<blocks,threads>>>调用的范围相同，但由于分支只有0号和1号线程执行ptr[threadIdx.x]++。这正确吗？ - scatman