NVCC和NVRTC是否支持尾部呼叫优化？

Question

我在F＃中做了一个小巧的函数式语言，它将编译为C++（大约第四次），并且想知道Nvidia的编译器现在是否具有此功能。我希望他们这样做，因为它可以节省我不得不实施元组的努力，而谷歌搜索什么也没有。

但是，在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 -02 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: 

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