Home
last modified time | relevance | path

Searched refs:src0 (Results 1 – 25 of 141) sorted by relevance

123456

/llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/
H A Dmad-mix.ll12 %src0.ext = fpext half %src0 to float
24 %src0.hi = lshr i32 %src0, 16
27 %src0.i16 = trunc i32 %src0.hi to i16
49 %src0.ext = fpext half %src0.hi to float
109 %src0.ext = fpext half %src0 to float
112 %src0.ext.neg = fneg float %src0.ext
123 %src0.ext = fpext half %src0 to float
142 %src0.ext = fpext half %src0 to float
159 %src0.ext = fpext half %src0 to float
462 %src0.neg = fsub half -0.0, %src0
[all …]
H A Dllvm.amdgcn.permlane.ll13 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 0, i1 0)
22 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 1, i32 2, i1 0, i1 0)
34 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 4660, i32 49617, i1 0, i1 0)
50 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %tidx, i32 %tidy, i1 0, i1 0)
62 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %tidx, i32 %src2, i1 0, i1 0)
74 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %tidy, i1 0, i1 0)
83 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 1, i1 0)
92 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 0, i1 1)
101 %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 1, i1 1)
119 %v = call i32 @llvm.amdgcn.permlanex16(i32 %src0, i32 %src0, i32 1, i32 2, i1 0, i1 0)
[all …]
H A Dllvm.amdgcn.fmed3.f16.ll5 define amdgpu_kernel void @test_fmed3_f16(half addrspace(1)* %out, i32 %src0.arg, i32 %src1.arg, i3…
6 %src0.f16 = trunc i32 %src0.arg to i16
7 %src0 = bitcast i16 %src0.f16 to half
12 %mad = call half @llvm.amdgcn.fmed3.f16(half %src0, half %src1, half %src2)
19 define amdgpu_kernel void @test_fmed3_srcmods_f16(half addrspace(1)* %out, i32 %src0.arg, i32 %src1…
20 %src0.f16 = trunc i32 %src0.arg to i16
21 %src0 = bitcast i16 %src0.f16 to half
26 %src0.fneg = fsub half -0.0, %src0
30 %mad = call half @llvm.amdgcn.fmed3.f16(half %src0.fneg, half %src1.fabs, half %src2.fneg.fabs)
H A Dmad-mix-hi.ll10 %src0.ext = fpext half %src0 to float
13 %result = tail call float @llvm.fmuladd.f32(float %src0.ext, float %src1.ext, float %src2.ext)
26 %src0.ext = fpext half %src0 to float
29 %result = tail call float @llvm.fmuladd.f32(float %src0.ext, float %src1.ext, float %src2.ext)
41 %src0.ext = fpext half %src0 to float
56 define i32 @v_mad_mixhi_f16_f16lo_f16lo_f16lo_intpack(half %src0, half %src1, half %src2) #0 {
57 %src0.ext = fpext half %src0 to float
74 %src0.ext = fpext half %src0 to float
91 %src0.ext = fpext half %src0 to float
107 %src0.ext = fpext half %src0 to float
[all …]
H A Dllvm.amdgcn.fmed3.ll6 define amdgpu_kernel void @test_fmed3(float addrspace(1)* %out, float %src0, float %src1, float %sr…
7 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float %src2)
15 %src0.fneg = fsub float -0.0, %src0
27 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float %src2)
37 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float %src2)
49 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float %src2)
59 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float 0.0)
71 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float %src1, float 0.0)
82 define amdgpu_kernel void @test_fneg_fmed3_r_inv2pi_0(float addrspace(1)* %out, float %src0) #1 {
83 %med3 = call float @llvm.amdgcn.fmed3.f32(float %src0, float 0x3FC45F3060000000, float 0.0)
[all …]
H A Dmad-mix-lo.ll26 %src0.ext = fpext half %src0 to float
42 %src0.ext = fpext half %src0 to float
57 %src0.ext = fpext half %src0 to float
75 %src0.ext = fpext half %src0 to float
98 %src0.ext = fpext <2 x half> %src0 to <2 x float>
119 %src0.ext = fpext <3 x half> %src0 to <3 x float>
143 %src0.ext = fpext <4 x half> %src0 to <4 x float>
162 %src0.ext = fpext <2 x half> %src0 to <2 x float>
188 %src0.ext = fpext <3 x half> %src0 to <3 x float>
216 %src0.ext = fpext <4 x half> %src0 to <4 x float>
[all …]
/llvm-project-15.0.7/clang/lib/Headers/
H A Damxintrin.h152 __builtin_ia32_tdpbssd((dst), (src0), (src1))
171 __builtin_ia32_tdpbsud((dst), (src0), (src1))
190 __builtin_ia32_tdpbusd((dst), (src0), (src1))
209 __builtin_ia32_tdpbuud((dst), (src0), (src1))
227 __builtin_ia32_tdpbf16ps((dst), (src0), (src1))
362 dst->tile = _tile_dpbssd_internal(src0.row, src1.col, src0.col, dst->tile, in __tile_dpbssd()
363 src0.tile, src1.tile); in __tile_dpbssd()
385 dst->tile = _tile_dpbsud_internal(src0.row, src1.col, src0.col, dst->tile, in __tile_dpbsud()
408 dst->tile = _tile_dpbusd_internal(src0.row, src1.col, src0.col, dst->tile, in __tile_dpbusd()
431 dst->tile = _tile_dpbuud_internal(src0.row, src1.col, src0.col, dst->tile, in __tile_dpbuud()
[all …]
/llvm-project-15.0.7/llvm/test/TableGen/
H A DDefaultOpsGlobalISel.td114 // CHECK: // (fround:{ *:[f32] } f32:{ *:[f32] }:$src0) => (FBAR:{ *:[f32] } f32:{ *:[f32] }:$src…
122 // CHECK: // (ftrunc:{ *:[f32] } f32:{ *:[f32] }:$src0) => (FFOO:{ *:[f32] } FPR32:{ *:[f32] }:$s…
138 (ftrunc f32:$src0),
139 (FFOO FPR32:$src0)
145 [(set FPR32:$dst, (fround f32:$src0))]>;
157 (fcos (SelectOMod f32:$src0, i32:$omod)),
158 (FLAMP FPR32:$src0, omod:$omod)
163 (fsqrt (SelectClamp f32:$src0, i1:$clamp)),
164 (FLAMP FPR32:$src0, 93, clamp:$clamp)
172 (fexp2 (SelectClamp f32:$src0, i1:$clamp)),
[all …]
H A DGlobalISelEmitter-SDNodeXForm-timm.td18 def SLEEP : I<(outs), (ins i32imm:$src0), []>;
19 def FOO : I<(outs GPR32:$dst), (ins GPR32:$src0, i32imm:$src1), []>;
26 (int_mytarget_foo i32:$src0, (i32 timm:$src1)),
27 (FOO GPR32:$src0, (shiftl_1 $src1))
33 …enderer, /*InsnID*/0, /*OldInsnID*/0, /*OpIdx*/1, /*OperandRenderer*/GICR_renderShiftImml1, // src0
35 (int_mytarget_sleep (i32 timm:$src0)),
36 (SLEEP (shiftl_1 $src0))
/llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/GlobalISel/
H A Dinst-select-abs.mir20 ; GFX6-NEXT: %src0:sreg_32 = COPY $sgpr0
29 %src0:sgpr(s32) = COPY $sgpr0
31 %ineg:sgpr(s32) = G_SUB %zero, %src0
32 %smax:sgpr(s32) = G_SMAX %src0, %ineg
58 %src0:sgpr(s32) = COPY $sgpr0
60 %ineg:sgpr(s32) = G_SUB %zero, %src0
61 %smax:sgpr(s32) = G_SMAX %ineg, %src0
91 %src0:vgpr(s32) = COPY $vgpr0
93 %ineg:vgpr(s32) = G_SUB %zero, %src0
94 %smax:vgpr(s32) = G_SMAX %src0, %ineg
[all …]
H A DshlN_add.ll25 %shl = shl i32 %src0, 1
46 %shl = shl i32 %src0, 2
67 %shl = shl i32 %src0, 3
88 %shl = shl i32 %src0, 4
99 %shl = shl i32 %src0, 5
124 %shl = shl i32 %src0, 1
149 %shl = shl i32 %src0, 2
174 %shl = shl i32 %src0, 3
199 %shl = shl i32 %src0, 4
224 %shl = shl i32 %src0, 5
[all …]
H A Dorn2.ll23 %or = or i32 %src0, %not.src1
43 %or = or i32 %not.src1, %src0
66 %or = or i32 %src0, %not.src1
91 %or0 = or i32 %src0, %not.src2
114 %or = or i32 %src0, %not.src1
131 %or = or i32 %src0, %not.src1
149 %or = or i32 %src0, %not.src1
170 %or = or i64 %src0, %not.src1
190 %or = or i64 %not.src1, %src0
241 %or = or i64 %src0, %not.src1
[all …]
H A Dandn2.ll23 %and = and i32 %src0, %not.src1
43 %and = and i32 %not.src1, %src0
66 %and = and i32 %src0, %not.src1
91 %and0 = and i32 %src0, %not.src2
114 %and = and i32 %src0, %not.src1
131 %and = and i32 %src0, %not.src1
149 %and = and i32 %src0, %not.src1
170 %and = and i64 %src0, %not.src1
190 %and = and i64 %not.src1, %src0
241 %and = and i64 %src0, %not.src1
[all …]
/llvm-project-15.0.7/llvm/lib/Target/AMDGPU/
H A DAMDGPUInstrInfo.td185 // out = (src0 + src1 > 0xFFFFFFFF) ? 1 : 0
188 // out = (src1 > src0) ? 1 : 0
249 // src0: vec4(src, 0, 0, mask)
325 // i32 or f32 src0
386 [(int_amdgcn_ldexp node:$src0, node:$src1),
387 (AMDGPUldexp_impl node:$src0, node:$src1)]>;
390 [(int_amdgcn_class node:$src0, node:$src1),
391 (AMDGPUfp_class_impl node:$src0, node:$src1)]>;
414 [(int_amdgcn_cvt_pkrtz node:$src0, node:$src1),
438 [(int_amdgcn_mul_u24 node:$src0, node:$src1),
[all …]
H A DAMDGPUInstructions.td172 (ops node:$src0),
173 (op $src0),
183 (op $src0, $src1),
258 def csh_mask_16 : PatFrag<(ops node:$src0), (and node:$src0, imm),
263 def csh_mask_32 : PatFrag<(ops node:$src0), (and node:$src0, imm),
268 def csh_mask_64 : PatFrag<(ops node:$src0), (and node:$src0, imm),
299 (ops node:$src0), (srl_oneuse node:$src0, (i32 16))
304 (ops node:$src0), (i16 (trunc (i32 (srl_16 node:$src0))))
694 (dt rc:$src0)
707 (BIT_ALIGN $src0, $src0, $src1)
[all …]
H A DSIInstructions.td71 let DisableEncoding = "$src0", Constraints = "$src0 = $vdst" in {
81 } // End DisableEncoding = "$src0", Constraints = "$src0 = $vdst"
199 (ins VGPR_32:$src0),
203 (ins VGPR_32:$src0),
2208 (S_NOT_B64 $src0)
2213 (S_NOT_B64 $src0)
2247 (S_NOT_B32 $src0)
2252 (S_NOT_B32 $src0)
2669 (COPY $src0)
2758 (COPY $src0)
[all …]
/llvm-project-15.0.7/llvm/docs/AMDGPU/
H A DAMDGPUAsmGFX7.rst1154 …v_cmp_eq_f32 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1155 …v_cmp_eq_f64 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1156 …v_cmp_eq_i32 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1157 …v_cmp_eq_i64 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1158 …v_cmp_eq_u32 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1159 …v_cmp_eq_u64 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1160 …v_cmp_f_f32 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1161 …v_cmp_f_f64 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1162 …v_cmp_f_i32 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1163 …v_cmp_f_i64 :ref:`vcc<amdgpu_synid_gfx7_vcc>`, :ref:`src0<amdgpu_synid_gfx…
[all …]
H A DAMDGPUAsmGFX90a.rst1661 …v_cmp_eq_f16 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1663 …v_cmp_eq_f32 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1665 …v_cmp_eq_f64 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1666 …v_cmp_eq_i16 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1668 …v_cmp_eq_i32 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1670 …v_cmp_eq_i64 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1671 …v_cmp_eq_u16 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1673 …v_cmp_eq_u32 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1675 …v_cmp_eq_u64 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
1676 …v_cmp_f_f16 :ref:`vcc<amdgpu_synid_gfx90a_vcc>`, :ref:`src0<amdgpu_synid_g…
[all …]
H A DAMDGPUAsmGFX9.rst1698 …v_cmp_eq_f16 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1700 …v_cmp_eq_f32 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1702 …v_cmp_eq_f64 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1703 …v_cmp_eq_i16 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1705 …v_cmp_eq_i32 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1707 …v_cmp_eq_i64 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1708 …v_cmp_eq_u16 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1710 …v_cmp_eq_u32 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1712 …v_cmp_eq_u64 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1713 …v_cmp_f_f16 :ref:`vcc<amdgpu_synid_gfx9_vcc>`, :ref:`src0<amdgpu_synid_gfx…
[all …]
H A DAMDGPUAsmGFX940.rst1680 …v_cmp_eq_f16 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1682 …v_cmp_eq_f32 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1684 …v_cmp_eq_f64 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1685 …v_cmp_eq_i16 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1687 …v_cmp_eq_i32 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1689 …v_cmp_eq_i64 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1690 …v_cmp_eq_u16 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1692 …v_cmp_eq_u32 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1694 …v_cmp_eq_u64 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
1695 …v_cmp_f_f16 :ref:`vcc<amdgpu_synid_gfx940_vcc>`, :ref:`src0<amdgpu_synid_g…
[all …]
H A DAMDGPUAsmGFX1030.rst1974 …v_cmpx_eq_f16 :ref:`src0<amdgpu_synid_gfx1030_src_5bb6f2>`, :ref:`v…
1975 …v_cmpx_eq_f32 :ref:`src0<amdgpu_synid_gfx1030_src_5bb6f2>`, :ref:`v…
1976 …v_cmpx_eq_f64 :ref:`src0<amdgpu_synid_gfx1030_src_b9c62f>`, :ref:`v…
1977 …v_cmpx_eq_i16 :ref:`src0<amdgpu_synid_gfx1030_src_364d7c>`, :ref:`v…
1978 …v_cmpx_eq_i32 :ref:`src0<amdgpu_synid_gfx1030_src_5bb6f2>`, :ref:`v…
1979 …v_cmpx_eq_i64 :ref:`src0<amdgpu_synid_gfx1030_src_b9c62f>`, :ref:`v…
1980 …v_cmpx_eq_u16 :ref:`src0<amdgpu_synid_gfx1030_src_364d7c>`, :ref:`v…
1981 …v_cmpx_eq_u32 :ref:`src0<amdgpu_synid_gfx1030_src_5bb6f2>`, :ref:`v…
1982 …v_cmpx_eq_u64 :ref:`src0<amdgpu_synid_gfx1030_src_b9c62f>`, :ref:`v…
1983 …v_cmpx_f_f16 :ref:`src0<amdgpu_synid_gfx1030_src_5bb6f2>`, :ref:`v…
[all …]
H A DAMDGPUAsmGFX10.rst2061 …v_cmpx_eq_f16 :ref:`src0<amdgpu_synid_gfx10_src_823582>`, :ref:`vsr…
2062 …v_cmpx_eq_f32 :ref:`src0<amdgpu_synid_gfx10_src_823582>`, :ref:`vsr…
2063 …v_cmpx_eq_f64 :ref:`src0<amdgpu_synid_gfx10_src_d5cd94>`, :ref:`vsr…
2064 …v_cmpx_eq_i16 :ref:`src0<amdgpu_synid_gfx10_src_e9e6db>`, :ref:`vsr…
2065 …v_cmpx_eq_i32 :ref:`src0<amdgpu_synid_gfx10_src_823582>`, :ref:`vsr…
2066 …v_cmpx_eq_i64 :ref:`src0<amdgpu_synid_gfx10_src_d5cd94>`, :ref:`vsr…
2067 …v_cmpx_eq_u16 :ref:`src0<amdgpu_synid_gfx10_src_e9e6db>`, :ref:`vsr…
2068 …v_cmpx_eq_u32 :ref:`src0<amdgpu_synid_gfx10_src_823582>`, :ref:`vsr…
2069 …v_cmpx_eq_u64 :ref:`src0<amdgpu_synid_gfx10_src_d5cd94>`, :ref:`vsr…
2070 …v_cmpx_f_f16 :ref:`src0<amdgpu_synid_gfx10_src_823582>`, :ref:`vsr…
[all …]
H A DAMDGPUAsmGFX8.rst1439 …v_cmp_eq_f16 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1441 …v_cmp_eq_f32 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1443 …v_cmp_eq_f64 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1444 …v_cmp_eq_i16 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1446 …v_cmp_eq_i32 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1448 …v_cmp_eq_i64 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1449 …v_cmp_eq_u16 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1451 …v_cmp_eq_u32 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1453 …v_cmp_eq_u64 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
1454 …v_cmp_f_f16 :ref:`vcc<amdgpu_synid_gfx8_vcc>`, :ref:`src0<amdgpu_synid_gfx…
[all …]
H A DAMDGPUAsmGFX906.rst43 …v_fmac_f32 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_…
45 …v_xnor_b32 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_…
47 …v_xnor_b32_sdwa :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_…
56 …v_fmac_f32_e64 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_…
57 …v_xnor_b32_e64 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_…
67 …v_dot2_i32_i16 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_gfx906_sr…
68 …v_dot2_u32_u16 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_gfx906_sr…
69 …v_dot4_i32_i8 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_gfx906_sr…
70 …v_dot4_u32_u8 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_gfx906_sr…
71 …v_dot8_i32_i4 :ref:`vdst<amdgpu_synid_gfx906_vdst>`, :ref:`src0<amdgpu_synid_gfx906_sr…
[all …]
/llvm-project-15.0.7/clang/test/CodeGenOpenCL/
H A Dbuiltins-amdgcn.cl654 *out = __builtin_amdgcn_mbcnt_lo(src0, src1);
660 *out = __builtin_amdgcn_mbcnt_hi(src0, src1);
666 *out = __builtin_amdgcn_alignbit(src0, src1, src2);
678 *out = __builtin_amdgcn_ubfe(src0, src1, src2);
684 *out = __builtin_amdgcn_sbfe(src0, src1, src2);
690 *out = __builtin_amdgcn_cvt_pkrtz(src0, src1);
708 *out = __builtin_amdgcn_cvt_pk_i16(src0, src1);
714 *out = __builtin_amdgcn_cvt_pk_u16(src0, src1);
726 *out = __builtin_amdgcn_sad_u8(src0, src1, src2);
732 *out = __builtin_amdgcn_msad_u8(src0, src1, src2);
[all …]

123456