NVCC 和 NVRTC 是否支持尾调用优化？

Question

我正在用 F# 编写一种小型函数式语言，它将编译为 C++（大约第四次），我想知道 Nvidia 的编译器现在是否具有此功能。我希望他们这样做，因为这样可以节省我实现元组的工作，而且 Google 搜索什么也找不到。

但是，在 Cuda 用户指南中也没有搜索 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;
}

来自 cubin 的 nvdisasm:

//--------------------- .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:

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

NVCC 和 NVRTC 是否支持尾调用优化？

Do NVCC and NVRTC support tail call optimization?

c++

recursion

cuda

tail-recursion