@@ -7,7 +7,7 @@ target triple = "nvptx-nvidia-cuda"
77declare i32 @llvm.nvvm.fshr.clamp.i32 (i32 , i32 , i32 )
88declare i32 @llvm.nvvm.fshl.clamp.i32 (i32 , i32 , i32 )
99
10- define i32 @fshr_clamp_r (i32 %a , i32 %b , i32 %c ) {
10+ define i32 @fshr_clamp_r (i32 %hi , i32 %lo , i32 %n ) {
1111; CHECK-LABEL: fshr_clamp_r(
1212; CHECK: {
1313; CHECK-NEXT: .reg .b32 %r<5>;
@@ -19,11 +19,11 @@ define i32 @fshr_clamp_r(i32 %a, i32 %b, i32 %c) {
1919; CHECK-NEXT: shf.r.clamp.b32 %r4, %r2, %r1, %r3;
2020; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
2121; CHECK-NEXT: ret;
22- %call = call i32 @llvm.nvvm.fshr.clamp.i32 (i32 %a , i32 %b , i32 %c )
22+ %call = call i32 @llvm.nvvm.fshr.clamp.i32 (i32 %hi , i32 %lo , i32 %n )
2323 ret i32 %call
2424}
2525
26- define i32 @fshl_clamp_r (i32 %a , i32 %b , i32 %c ) {
26+ define i32 @fshl_clamp_r (i32 %hi , i32 %lo , i32 %n ) {
2727; CHECK-LABEL: fshl_clamp_r(
2828; CHECK: {
2929; CHECK-NEXT: .reg .b32 %r<5>;
@@ -35,11 +35,11 @@ define i32 @fshl_clamp_r(i32 %a, i32 %b, i32 %c) {
3535; CHECK-NEXT: shf.l.clamp.b32 %r4, %r2, %r1, %r3;
3636; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
3737; CHECK-NEXT: ret;
38- %call = call i32 @llvm.nvvm.fshl.clamp.i32 (i32 %a , i32 %b , i32 %c )
38+ %call = call i32 @llvm.nvvm.fshl.clamp.i32 (i32 %hi , i32 %lo , i32 %n )
3939 ret i32 %call
4040}
4141
42- define i32 @fshr_clamp_i (i32 %a , i32 %b , i32 %c ) {
42+ define i32 @fshr_clamp_i (i32 %hi , i32 %lo ) {
4343; CHECK-LABEL: fshr_clamp_i(
4444; CHECK: {
4545; CHECK-NEXT: .reg .b32 %r<4>;
@@ -50,11 +50,11 @@ define i32 @fshr_clamp_i(i32 %a, i32 %b, i32 %c) {
5050; CHECK-NEXT: shf.r.clamp.b32 %r3, %r2, %r1, 3;
5151; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
5252; CHECK-NEXT: ret;
53- %call = call i32 @llvm.nvvm.fshr.clamp.i32 (i32 %a , i32 %b , i32 3 )
53+ %call = call i32 @llvm.nvvm.fshr.clamp.i32 (i32 %hi , i32 %lo , i32 3 )
5454 ret i32 %call
5555}
5656
57- define i32 @fshl_clamp_i (i32 %a , i32 %b , i32 %c ) {
57+ define i32 @fshl_clamp_i (i32 %hi , i32 %lo ) {
5858; CHECK-LABEL: fshl_clamp_i(
5959; CHECK: {
6060; CHECK-NEXT: .reg .b32 %r<4>;
@@ -65,6 +65,6 @@ define i32 @fshl_clamp_i(i32 %a, i32 %b, i32 %c) {
6565; CHECK-NEXT: shf.l.clamp.b32 %r3, %r2, %r1, 3;
6666; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
6767; CHECK-NEXT: ret;
68- %call = call i32 @llvm.nvvm.fshl.clamp.i32 (i32 %a , i32 %b , i32 3 )
68+ %call = call i32 @llvm.nvvm.fshl.clamp.i32 (i32 %hi , i32 %lo , i32 3 )
6969 ret i32 %call
7070}
0 commit comments