如果所有的翘曲都走一条路径，CUDA中的分支可以被忽略吗?如果是这样，有一种方法，我可以给编译器/运行时这个信息?<

Question

假设我们有如下代码(我没有编译这个，它可能是错误的)

__global__ void myKernel()
{
int data = someArray[threadIdx.x];

if (data == 0) {
funcA();
} else {
funcB();
}
}

现在假设有1024个线程块正在运行，并且someArray都是零。

进一步假设funcB()运行成本高，而funcA()运行成本低。

我假设编译器必须按顺序发出两条路径，比如先执行funcA，然后执行funcB。这不是理想。

有没有办法提示CUDA不要这样做?还是运行时注意"没有线程是活动所以我将跳过所有的指令作为我看到them" ?

或者更好，如果分支是这样的(再次强调，没有编译这个，但它说明了我想要传达的内容)

__constant__ int constantNumber;
__global__ void myKernel()
{
if (constantNumber == 123) {
funcA();
} else {
funcB();
}
}

，然后在启动内核之前将constantNumber设置为123。这还会导致两条路都走吗?

Answer 1

这可以使用__builtin_assume来实现。https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html __builtin_assume

引用文档:

void __builtin_assume(bool exp)允许编译器假定布尔参数为真。如果参数在运行时不为真，则行为未定义。参数不会被求值，因此任何副作用都将被忽略。