Home
last modified time | relevance | path

Searched refs:b64 (Results 1 – 25 of 108) sorted by relevance

12345

/llvm-project-15.0.7/llvm/test/CodeGen/NVPTX/
H A Dmbarrier.ll11 ; CHECK_PTX32: mbarrier.init.b64 [%r{{[0-9]+}}], %r{{[0-9]+}};
12 ; CHECK_PTX64: mbarrier.init.b64 [%rd{{[0-9]+}}], %r{{[0-9]+}};
19 ; CHECK_PTX32: mbarrier.init.shared.b64 [%r{{[0-9]+}}], %r{{[0-9]+}};
20 ; CHECK_PTX64: mbarrier.init.shared.b64 [%rd{{[0-9]+}}], %r{{[0-9]+}};
30 ; CHECK_PTX32: mbarrier.inval.b64 [%r{{[0-1]+}}];
31 ; CHECK_PTX64: mbarrier.inval.b64 [%rd{{[0-1]+}}];
38 ; CHECK_PTX32: mbarrier.inval.shared.b64 [%r{{[0-1]+}}];
39 ; CHECK_PTX64: mbarrier.inval.shared.b64 [%rd{{[0-1]+}}];
49 ; CHECK_PTX32: mbarrier.arrive.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}];
50 ; CHECK_PTX64: mbarrier.arrive.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}];
[all …]
H A Drotate_64.ll4 declare i64 @llvm.nvvm.rotate.b64(i64, i32)
5 declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
9 ; CHECK: shl.b64 [[LHS:%.*]], [[RD1:%.*]], 3;
10 ; CHECK: shr.b64 [[RHS:%.*]], [[RD1]], 61;
13 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 3)
19 ; CHECK: shl.b64 [[LHS:%.*]], [[RD1:%.*]], 61;
20 ; CHECK: shr.b64 [[RHS:%.*]], [[RD1]], 3;
23 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 3)
H A Dshift-parts.ll6 ; CHECK: shl.b64
10 ; CHECK: or.b64
12 ; CHECK: shl.b64
14 ; CHECK: selp.b64
15 ; CHECK: shl.b64
27 ; CHECK: shl.b64
28 ; CHECK: or.b64
32 ; CHECK: selp.b64
H A Drotate.ll8 declare i64 @llvm.nvvm.rotate.b64(i64, i32)
9 declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
26 ; SM20: shl.b64
28 ; SM20: shr.b64
32 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 %b)
39 ; SM20: shr.b64
41 ; SM20: shl.b64
45 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 %b)
H A Datomicrmw-expand.ll9 ; SM30: atom.cas.b64
40 ; SM30: atom.cas.b64
41 ; SM60: atom.and.b64
43 ; SM30: atom.cas.b64
44 ; SM60: atom.or.b64
46 ; SM30: atom.cas.b64
47 ; SM60: atom.xor.b64
49 ; SM30: atom.cas.b64
72 ; SM30: atom.cas.b64
75 ; SM30: atom.cas.b64
[all …]
H A Dparam-load-store.ll696 ; CHECK: .param .b64 param0;
698 ; CHECK: .param .b64 retval0;
714 ; CHECK: .param .b64 param0;
716 ; CHECK: .param .b64 retval0;
732 ; CHECK: .param .b64 param0;
734 ; CHECK: .param .b64 retval0;
751 ; CHECK: .param .b64 param0;
753 ; CHECK: .param .b64 retval0;
770 ; CHECK: .param .b64 param0;
787 ; CHECK: .param .b64 param0;
[all …]
H A Dasync-copy.ll44 ; CHECK_PTX32: cp.async.mbarrier.arrive.b64 [%r{{[0-9]+}}];
45 ; CHECK_PTX64: cp.async.mbarrier.arrive.b64 [%rd{{[0-9]+}}];
52 ; CHECK_PTX32: cp.async.mbarrier.arrive.shared.b64 [%r{{[0-9]+}}];
53 ; CHECK_PTX64: cp.async.mbarrier.arrive.shared.b64 [%rd{{[0-9]+}}];
60 ; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.b64 [%r{{[0-9]+}}];
61 ; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.b64 [%rd{{[0-9]+}}];
68 ; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.shared.b64 [%r{{[0-9]+}}];
69 ; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.shared.b64 [%rd{{[0-9]+}}];
H A Dproxy-reg-erasure-ptx.ll64 ; PTX-DAG: ld.param.b64 [[LD:%rd[0-9]+]], [retval0+0];
67 ; PTX-WITHOUT-DAG: mov.b64 [[PROXY:%rd[0-9]+]], [[LD]];
68 ; PTX-WITHOUT-DAG: st.param.b64 [func_retval0+0], [[PROXY]];
69 ; PTX-WITH-DAG: st.param.b64 [func_retval0+0], [[LD]];
79 ; PTX-DAG: ld.param.v2.b64 {[[LD0:%rd[0-9]+]], [[LD1:%rd[0-9]+]]}, [retval0+0];
82 ; PTX-WITHOUT-DAG: mov.b64 [[PROXY0:%rd[0-9]+]], [[LD0]];
83 ; PTX-WITHOUT-DAG: mov.b64 [[PROXY1:%rd[0-9]+]], [[LD1]];
84 ; PTX-WITHOUT-DAG: st.param.v2.b64 [func_retval0+0], {[[PROXY0]], [[PROXY1]]};
85 ; PTX-WITH-DAG: st.param.v2.b64 [func_retval0+0], {[[LD0]], [[LD1]]};
H A Dmatch.ll31 ; CHECK: match.any.sync.b64 [[V0:%r[0-9]+]], [[VALUE]], [[MASK]];
33 ; CHECK: match.any.sync.b64 [[V1:%r[0-9]+]], [[VALUE]], 1;
35 ; CHECK: match.any.sync.b64 [[V2:%r[0-9]+]], 2, [[MASK]];
37 ; CHECK: match.any.sync.b64 [[V3:%r[0-9]+]], 4, 3;
89 ; CHECK: match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
94 ; CHECK: match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
99 ; CHECK: match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], 2;
104 ; CHECK: match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, 4, 3;
H A Dcall-with-alloca-buffer.ll51 ; CHECK: .param .b64 param0;
52 ; CHECK-NEXT: st.param.b64 [param0+0], %rd[[A_REG]]
53 ; CHECK-NEXT: .param .b64 param1;
54 ; CHECK-NEXT: st.param.b64 [param1+0], %rd[[SP_REG]]
H A Di128-param.ll33 ; CHECK-NEXT: st.param.v2.b64 [param0+0], {%[[REG0]], %[[REG1]]}
35 ; CHECK-NEXT: st.param.v2.b64 [param1+0], {%[[REG2]], %[[REG3]]}
52 ; CHECK: st.param.v2.b64 [param0+0], {%[[REG0]], %[[REG1]]}
54 ; CHECK: st.param.v2.b64 [param1+0], {%[[REG2]], %[[REG3]]}
H A Dctlz.ll32 ; PTX's clz.b64 returns a 32-bit value, but LLVM's intrinsic returns a 64-bit
37 ; CHECK-NEXT: clz.b64
47 ; CHECK-NEXT: clz.b64
56 ; natural return width of ptx's clz.b64 instruction. No conversions should be
61 ; CHECK-NEXT: clz.b64
71 ; CHECK-NEXT: clz.b64
H A Di128-struct.ll11 ; CHECK: st.param.v2.b64 [func_retval0+0], {%[[REG1:rd[0-9]+]], %[[REG2:rd[0-9]+]]};
12 ; CHECK: st.param.v2.b64 [func_retval0+16], {%[[REG3:rd[0-9]+]], %[[REG4:rd[0-9]+]]};
H A Dvec-param-load.ll63 ; CHECK-DAG: st.param.v2.b64 [func_retval0+0], {[[V_0_1]]}
64 ; CHECK-DAG: st.param.v2.b64 [func_retval0+16], {[[V_2_3]]}
65 ; CHECK-DAG: st.param.v2.b64 [func_retval0+32], {[[V_4_5]]}
66 ; CHECK-DAG: st.param.v2.b64 [func_retval0+48], {[[V_6_7]]}
H A Di128-retval.ll7 ; CHECK: st.param.v2.b64 [func_retval0+0], {%[[REG0]], %[[REG1]]}
20 ; CHECK: ld.param.v2.b64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [retval0+0];
H A Dlibcall-fulfilled.ll21 ; CHECK-NEXT: ld.param.v2.b64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [retval0+0];
27 ; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%[[REG0]], %[[REG1]]};
H A Datomics-with-scope.ll32 ; CHECK: atom.cta.exch.b64
36 ; CHECK: atom.sys.exch.b64
69 ; CHECK: atom.cta.and.b64
73 ; CHECK: atom.sys.and.b64
78 ; CHECK: atom.cta.or.b64
82 ; CHECK: atom.sys.or.b64
87 ; CHECK: atom.cta.xor.b64
91 ; CHECK: atom.sys.xor.b64
96 ; CHECK: atom.cta.cas.b64
100 ; CHECK: atom.sys.cas.b64
H A Dintrinsics.ll43 ; CHECK: brev.b64
57 ; CHECK: popc.b64
63 ; NVPTX popc.b64 returns an i32 even though @llvm.ctpop.i64 returns an i64, so
68 ; CHECK: popc.b64
/llvm-project-15.0.7/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/
H A Dnvptx-basic.ll.expected13 ; CHECK-NEXT: .reg .b64 %rd<17>;
35 ; CHECK-NEXT: st.param.v2.b64 [param0+0], {%rd6, %rd5};
36 ; CHECK-NEXT: st.param.v2.b64 [param0+16], {%rd8, %rd7};
43 ; CHECK-NEXT: ld.param.v2.b64 {%rd9, %rd10}, [retval0+0];
44 ; CHECK-NEXT: ld.param.v2.b64 {%rd11, %rd12}, [retval0+16];
70 ; CHECK-NEXT: .reg .b64 %rd<5>;
75 ; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd1, %rd2};
76 ; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4};
/llvm-project-15.0.7/compiler-rt/test/builtins/Unit/
H A Datomic_test.c157 uint64_t a64, b64; variable
176 b64 = U64(value); in set_b_values()
405 if (a64 != U64(ONES) || b64 != U64(ONES)) in test_compare_exchanges()
409 if (a64 != U64(V + m1) || b64 != U64(ONES)) in test_compare_exchanges()
452 b64 = __atomic_fetch_add_8(&a64, U64(ONES), model); in test_fetch_op()
453 if (b64 != U64(V + m) || a64 != U64(V + m + ONES)) in test_fetch_op()
474 b64 = __atomic_fetch_sub_8(&a64, U64(ONES), model); in test_fetch_op()
475 if (b64 != U64(V + m) || a64 != U64(V + m - ONES)) in test_fetch_op()
497 if (b64 != U64(V + m) || a64 != U64(V + m)) in test_fetch_op()
518 b64 = __atomic_fetch_or_8(&a64, U64(ONES), model); in test_fetch_op()
[all …]
/llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/
H A Dreadcyclecounter.ll11 ; GCN-DAG: s_load_{{dwordx2|b64}}
23 ; GETREG: global_store_{{dwordx2|b64}} v{{.+}}, v[[[VCNT1]]:[[ZERO]]]
26 ; GETREG: global_store_{{dwordx2|b64}} v{{.+}}, v[[[VCNT2]]:[[ZERO]]]
41 ; GCN-DAG: s_load_{{dword|b32|b64}}
/llvm-project-15.0.7/llvm/test/DebugInfo/NVPTX/
H A Ddbg-declare-alloca.ll13 ; CHECK: .param .b64 param0;
14 ; CHECK: st.param.b64 [param0+0], %rd1;
153 ; CHECK-NEXT: .b64 $L__func_begin0 // DW_AT_low_pc
154 ; CHECK-NEXT: .b64 $L__func_end0 // DW_AT_high_pc
156 ; CHECK-NEXT: .b64 $L__func_begin0 // DW_AT_low_pc
157 ; CHECK-NEXT: .b64 $L__func_end0 // DW_AT_high_pc
184 ; CHECK-NEXT: .b64 __local_depot0
/llvm-project-15.0.7/llvm/test/Analysis/CostModel/X86/
H A Dshuffle-transpose.ll71 … of 2 for instruction: %V64 = shufflevector <2 x float> %a64, <2 x float> %b64, <2 x i32> <i32 0, …
78 … of 2 for instruction: %V64 = shufflevector <2 x float> %a64, <2 x float> %b64, <2 x i32> <i32 0, …
85 … of 2 for instruction: %V64 = shufflevector <2 x float> %a64, <2 x float> %b64, <2 x i32> <i32 0, …
98 %V64 = shufflevector <2 x float> %a64, <2 x float> %b64, <2 x i32> <i32 0, i32 2>
107 …cost of 2 for instruction: %V64 = shufflevector <2 x i32> %a64, <2 x i32> %b64, <2 x i32> <i32 0, …
114 …cost of 2 for instruction: %V64 = shufflevector <2 x i32> %a64, <2 x i32> %b64, <2 x i32> <i32 0, …
121 …cost of 2 for instruction: %V64 = shufflevector <2 x i32> %a64, <2 x i32> %b64, <2 x i32> <i32 0, …
128 …cost of 1 for instruction: %V64 = shufflevector <2 x i32> %a64, <2 x i32> %b64, <2 x i32> <i32 0, …
134 %V64 = shufflevector <2 x i32> %a64, <2 x i32> %b64, <2 x i32> <i32 0, i32 2>
207 %V64 = shufflevector <4 x i16> %a64, <4 x i16> %b64, <4 x i32> <i32 0, i32 4, i32 2, i32 6>
[all …]
/llvm-project-15.0.7/llvm/test/CodeGen/VE/Scalar/
H A Drotr.ll12 %b64 = zext i32 %b to i64
13 %a.lr = lshr i64 %a, %b64
H A Drotl.ll12 %b64 = zext i32 %b to i64
13 %a.sl = shl i64 %a, %b64

12345