NVCC和NVRTC支持尾部调用优化吗？

Question

我正在用F#做一种小型的函数式语言，它将编译成C++ (大约是第四次)，我想知道Nvidia的编译器现在是否有这个特性。我希望他们这样做可以省去我实现元组的努力，而谷歌搜索却一无所获。

然而，在库达用户指南中也没有搜索tail call，所以我想这是不太可能的。

Answer 1

简短回答:是的！

很长的答案:这个代码

__device__ int i;

__device__ int f2(int val)
{
        if((val % 6) == 0)
                return val;

        return f2(val + 1);
}

__global__ void f1(int x)
{
            i = f2(x);
}

int main()
{
            return 0;
}

编撰：

nvcc -keep -O2 bla.cu

生成: bla.ptx

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21124049
// Cuda compilation tools, release 8.0, V8.0.44
// Based on LLVM 3.4svn
//

.version 5.0
.target sm_20
.address_size 64

        // .globl       _Z2f1i
.global .align 4 .u32 i;

.visible .entry _Z2f1i(
        .param .u32 _Z2f1i_param_0
)
{
        .reg .pred      %p<2>;
        .reg .b32       %r<10>;


        ld.param.u32    %r9, [_Z2f1i_param_0];

BB0_1:
        mov.u32         %r1, %r9;
        mul.hi.s32      %r4, %r1, 715827883;
        shr.u32         %r5, %r4, 31;
        add.s32         %r6, %r4, %r5;
        mul.lo.s32      %r7, %r6, 6;
        sub.s32         %r8, %r1, %r7;
        add.s32         %r9, %r1, 1;
        setp.ne.s32     %p1, %r8, 0;
        @%p1 bra        BB0_1;

        st.global.u32   [i], %r1;
        ret;
}

来自nvdisasm的cubin：

//--------------------- .text._Z2f1i              --------------------------
        .section        .text._Z2f1i,"ax",@progbits
        .sectioninfo    @"SHI_REGISTERS=6"
        .align  4
        .global         _Z2f1i
        .type           _Z2f1i,@function
        .size           _Z2f1i,(.L_13 - _Z2f1i)
        .other          _Z2f1i,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z2f1i:
.text._Z2f1i:
        /*0000*/         MOV R1, c[0x1][0x100];
        /*0008*/         MOV R0, c[0x0][0x20];
        /*0010*/         NOP;
        /*0018*/         NOP;
        /*0020*/         NOP;
        /*0028*/         NOP;
        /*0030*/         NOP;
        /*0038*/         NOP;
.L_2:
        /*0040*/         IMUL.HI R2, R0, c[0x10][0x0];
        /*0048*/         IMAD.U32.U32.HI R2, R2, 0x2, R2;
        /*0050*/         IMAD R2, -R2, 0x6, R0;
        /*0058*/         ISETP.NE.AND P0, PT, R2, RZ, PT;
        /*0060*/         MOV R2, R0;
        /*0068*/         IADD R0, R0, 0x1;
        /*0070*/     @P0 BRA `(.L_2);
        /*0078*/         MOV R4, c[0xe][0x0];
        /*0080*/         MOV R5, c[0xe][0x4];
        /*0088*/         ST.E [R4], R2;
        /*0090*/         EXIT;
.L_13:

在这两种情况下，都有谓词分支，但没有调用。