| /llvm-project-15.0.7/llvm/test/CodeGen/AMDGPU/ |
| H A D | mad-mix.ll | 12 %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 D | llvm.amdgcn.permlane.ll | 13 %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 D | llvm.amdgcn.fmed3.f16.ll | 5 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 D | mad-mix-hi.ll | 10 %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 D | llvm.amdgcn.fmed3.ll | 6 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 D | mad-mix-lo.ll | 26 %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 D | amxintrin.h | 152 __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 D | DefaultOpsGlobalISel.td | 114 // 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 D | GlobalISelEmitter-SDNodeXForm-timm.td | 18 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 D | inst-select-abs.mir | 20 ; 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 D | shlN_add.ll | 25 %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 D | orn2.ll | 23 %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 D | andn2.ll | 23 %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 D | AMDGPUInstrInfo.td | 185 // 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 D | AMDGPUInstructions.td | 172 (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 D | SIInstructions.td | 71 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 D | AMDGPUAsmGFX7.rst | 1154 …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 D | AMDGPUAsmGFX90a.rst | 1661 …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 D | AMDGPUAsmGFX9.rst | 1698 …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 D | AMDGPUAsmGFX940.rst | 1680 …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 D | AMDGPUAsmGFX1030.rst | 1974 …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 D | AMDGPUAsmGFX10.rst | 2061 …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 D | AMDGPUAsmGFX8.rst | 1439 …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 D | AMDGPUAsmGFX906.rst | 43 …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 D | builtins-amdgcn.cl | 654 *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 …]
|