DO NVCC和NVRTC支持尾部呼叫优化
Do NVCC and NVRTC support tail call optimization?
我正在用f#进行小型功能语言,该语言将编译为C (大约第四次),并且想知道Nvidia的编译器现在是否具有此功能。我希望他们这样做,因为这可以节省我必须实施元组的努力,而Google搜索一无所获。
但是,在CUDA用户指南中搜索tail call
也不是,所以我想这不太可能。
简短答案:是!
长答案:此代码
__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:
在这两种情况下都有谓语分支,但没有调用。
相关文章:
- 呼叫运营商<<临时
- 如何建立使用模板函数的lambda函数的尾部返回类型
- 将尾部调用void(i32,..)位转换为llvm::函数以获取FnAttribute
- basic_string的前导/尾部不区分空格的特征
- 在带有尾部斜杠的路径上返回 std::filesystem::create_directories() 的值
- 编写一个函数来删除单链表中的节点(尾部除外),仅授予对该节点的访问权限
- 呼叫 QSound,它们之间有延迟 Qt C++
- 什么是呼叫说明
- 为什么make_shared在不同的呼叫中分配相同的地址?
- 错误:类型"std::__1::basic_string<char>"不提供呼叫运算符
- 设备IO控制呼叫崩溃
- 根据呼叫位置进行记忆
- 成员呼叫通过接线员<<
- 如何应用尾部调用优化
- QSerialPort 手动 RTS 开/关未同步呼叫
- 使用递归从尾部开始反转链表
- 将清除共享智能指针上的呼叫重置
- 来自 Pharo Smalltalk 的 ffi 呼叫上的分段错误
- C/C++ 在条件和阻止呼叫之间收到信号情报
- DO NVCC和NVRTC支持尾部呼叫优化