Home
last modified time | relevance | path

Searched refs:GFX908 (Results 1 – 25 of 27) sorted by relevance

12

/external/llvm-project/llvm/test/CodeGen/AMDGPU/
Dmai-inline.ll1 …llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX908 %s
4 ; GFX908: v_accvgpr_write [[AREG:a[0-9]+]], 1
5 ; GFX908: v_accvgpr_read [[VREG:v[0-9]+]], [[AREG]]
6 ; GFX908: global_store_dword v{{[0-9]+}}, [[VREG]], s{{\[[0-9]+:[0-9]+\]}}
17 ; GFX908: v_accvgpr_write_b32
18 ; GFX908: v_accvgpr_write_b32
19 ; GFX908: v_accvgpr_write_b32
20 ; GFX908: v_accvgpr_write_b32
21 ; GFX908: v_mfma_f32_4x4x1f32 a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9:]+}}]
22 ; GFX908: v_accvgpr_read_b32
[all …]
Dagpr-remat.ll2 …gcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GFX908 %s
7 ; GFX908-LABEL: remat_constant_voids_spill:
8 ; GFX908: ; %bb.0:
9 ; GFX908-NEXT: v_accvgpr_write_b32 a1, 1
10 ; GFX908-NEXT: v_accvgpr_write_b32 a5, 6
11 ; GFX908-NEXT: v_accvgpr_write_b32 a6, 7
12 ; GFX908-NEXT: v_accvgpr_write_b32 a7, 8
13 ; GFX908-NEXT: v_accvgpr_write_b32 a0, 9
14 ; GFX908-NEXT: v_accvgpr_write_b32 a2, 2
15 ; GFX908-NEXT: v_accvgpr_write_b32 a3, 3
[all …]
Dspill-vgpr-to-agpr.ll1 …llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX908 %s
7 ; GFX908-NOT: SCRATCH_RSRC
8 ; GFX908-DAG: v_accvgpr_write_b32 a0, v{{[0-9]}} ; Reload Reuse
9 ; GFX908-DAG: v_accvgpr_write_b32 a1, v{{[0-9]}} ; Reload Reuse
14 ; GFX908-NOT: buffer_
15 ; GFX908-DAG: v_accvgpr_read_b32 v{{[0-9]}}, a0 ; Reload Reuse
16 ; GFX908-DAG: v_accvgpr_read_b32 v{{[0-9]}}, a1 ; Reload Reuse
20 ; GFX908: ScratchSize: 0
60 ; GFX908-DAG: s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD0
61 ; GFX908-DAG: s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD1
[all …]
Delf-header-flags-sram-ecc.ll10 …dgcn -mcpu=gfx908 < %s | llvm-readobj -file-headers - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
38 ; SRAM-ECC-GFX908: Flags [ (0x230)
39 ; SRAM-ECC-GFX908: EF_AMDGPU_MACH_AMDGCN_GFX908 (0x30)
40 ; SRAM-ECC-GFX908: EF_AMDGPU_SRAM_ECC (0x200)
41 ; SRAM-ECC-GFX908: ]
Dglobal-atomics-fp.ll2 …-mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX908 %s
35 ; GFX908-NOT: v_add_f32
36 ; GFX908: global_atomic_add_f32 v{{[0-9]+}}, v{{[0-9]+}}, s
37 ; GFX908-NOT: s_cbranch_execnz
Dspill-agpr.ll1 …u=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX908,A2V %s
2 …o-agpr=0 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX908,A2M %s
8 ; GFX908-DAG: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a0 ; Reload Reuse
11 ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, v[[VSPILL]] ; Reload Reuse
37 ; GFX908-DAG: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a{{[0-9]+}} ; Reload Reuse
88 ; GFX908-DAG: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a0 ; Reload Reuse
91 ; GFX908: v_accvgpr_write_b32 a{{[0-9]+}}, v[[VSPILL]] ; Reload Reuse
Dllvm.amdgcn.sdot2.ll2 …lc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX908
10 ; GFX908: v_dot2_i32_i16 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
28 ; GFX908: v_dot2c_i32_i16_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
Dllvm.amdgcn.sdot8.ll2 …lc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck %s --check-prefixes=GCN,GFX908
12 ; GFX908: v_dot8_i32_i4 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} clamp{{$}}
32 ; GFX908: v_dot8c_i32_i4_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}{{$}}
Dspill-agpr-partially-undef.mir2 …achineinstrs -run-pass=prologepilog %s -o - | FileCheck -check-prefix=CHECK -check-prefix=GFX908 %s
Dhsa-note-no-func.ll26 …u=gfx908 --amdhsa-code-object-version=2 | FileCheck --check-prefix=HSA --check-prefix=HSA-GFX908 %s
52 ; HSA-GFX908: .hsa_code_object_isa 9,0,8,"AMD","AMDGPU"
Delf-header-flags-mach.ll53 …mcpu=gfx908 < %s | llvm-readobj -file-headers - | FileCheck --check-prefixes=ALL,ARCH-GCN,GFX908 %s
110 ; GFX908: EF_AMDGPU_MACH_AMDGCN_GFX908 (0x30)
111 ; GFX908-NEXT: EF_AMDGPU_SRAM_ECC (0x200)
Dglobal-saddr-atomics.gfx908.ll2 ; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx908 < %s | FileCheck -check-prefixes=GCN,GFX908 %s
/external/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/
Dllvm.amdgcn.global.atomic.fadd.ll2 …hsa -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GFX908 %s
5 ; GFX908-LABEL: global_atomic_fadd_f32:
6 ; GFX908: ; %bb.0:
7 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
8 ; GFX908-NEXT: global_atomic_add_f32 v[0:1], v2, off
9 ; GFX908-NEXT: s_waitcnt vmcnt(0)
10 ; GFX908-NEXT: s_setpc_b64 s[30:31]
16 ; GFX908-LABEL: global_atomic_fadd_f32_off_2048:
17 ; GFX908: ; %bb.0:
18 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
[all …]
Dirtranslator-call-implicit-args.ll3 …3d -mcpu=gfx908 -verify-machineinstrs -o - %s | FileCheck -enable-var-scope -check-prefix=GFX908 %s
61 ; GFX908-LABEL: name: test_call_external_void_func_i32
62 ; GFX908: bb.1 (%ir-block.1):
63 …; GFX908: liveins: $sgpr14, $sgpr15, $sgpr16, $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7…
64 ; GFX908: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr2
65 ; GFX908: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr1
66 ; GFX908: [[COPY2:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
67 ; GFX908: [[COPY3:%[0-9]+]]:sgpr_32 = COPY $sgpr16
68 ; GFX908: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr15
69 ; GFX908: [[COPY5:%[0-9]+]]:sgpr_32 = COPY $sgpr14
[all …]
Dllvm.amdgcn.sdot2.ll3 …sel -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX908 %s
14 ; GFX908-LABEL: v_sdot2:
15 ; GFX908: ; %bb.0:
16 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
17 ; GFX908-NEXT: v_dot2_i32_i16 v0, v0, v1, v2
18 ; GFX908-NEXT: s_setpc_b64 s[30:31]
37 ; GFX908-LABEL: v_sdot2_clamp:
38 ; GFX908: ; %bb.0:
39 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
40 ; GFX908-NEXT: v_dot2_i32_i16 v0, v0, v1, v2 clamp
[all …]
Dllvm.amdgcn.udot2.ll3 …sel -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX908 %s
14 ; GFX908-LABEL: v_udot2:
15 ; GFX908: ; %bb.0:
16 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
17 ; GFX908-NEXT: v_dot2_u32_u16 v0, v0, v1, v2
18 ; GFX908-NEXT: s_setpc_b64 s[30:31]
37 ; GFX908-LABEL: v_udot2_clamp:
38 ; GFX908: ; %bb.0:
39 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
40 ; GFX908-NEXT: v_dot2_u32_u16 v0, v0, v1, v2 clamp
[all …]
Dllvm.amdgcn.global.atomic.fadd-with-ret.ll1 …isel < %s -march=amdgcn -mcpu=gfx908 -verify-machineinstrs 2>&1 | FileCheck %s -check-prefix=GFX908
5 ; GFX908: error: {{.*}} return versions of fp atomics not supported
Dllvm.amdgcn.struct.buffer.atomic.fadd-with-ret.ll1 …isel < %s -march=amdgcn -mcpu=gfx908 -verify-machineinstrs 2>&1 | FileCheck %s -check-prefix=GFX908
5 ; GFX908: error: {{.*}} return versions of fp atomics not supported
/external/llvm-project/llvm/test/MC/Disassembler/AMDGPU/
Dmai.txt1 …lvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX908 %s
3 # GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
6 # GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
9 # GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]
12 # GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
15 # GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]
18 # GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
21 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03…
24 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x1…
27 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03…
[all …]
/external/llvm-project/llvm/test/Transforms/AtomicExpand/AMDGPU/
Dexpand-atomic-rmw-fadd.ll4 …pt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -atomic-expand %s | FileCheck -check-prefix=GFX908 %s
41 ; GFX908-LABEL: @test_atomicrmw_fadd_f32_flat(
42 ; GFX908-NEXT: [[TMP1:%.*]] = load float, float* [[PTR:%.*]], align 4
43 ; GFX908-NEXT: br label [[ATOMICRMW_START:%.*]]
44 ; GFX908: atomicrmw.start:
45 ; GFX908-NEXT: [[LOADED:%.*]] = phi float [ [[TMP1]], [[TMP0:%.*]] ], [ [[TMP6:%.*]], [[ATOMICRM…
46 ; GFX908-NEXT: [[NEW:%.*]] = fadd float [[LOADED]], [[VALUE:%.*]]
47 ; GFX908-NEXT: [[TMP2:%.*]] = bitcast float* [[PTR]] to i32*
48 ; GFX908-NEXT: [[TMP3:%.*]] = bitcast float [[NEW]] to i32
49 ; GFX908-NEXT: [[TMP4:%.*]] = bitcast float [[LOADED]] to i32
[all …]
/external/llvm-project/clang/include/clang/Basic/
DCuda.h76 GFX908, enumerator
/external/llvm-project/llvm/test/tools/llvm-readobj/ELF/
Damdgpu-elf-headers.test55 # RUN: yaml2obj %s -o %t -DCPU=GFX908
56 # RUN: llvm-readobj -h %t | FileCheck %s --match-full-lines -DFILE=%t -DCPU=GFX908 -DFLAGS=0x30
/external/llvm-project/clang/test/CodeGenOpenCL/
Damdgpu-features.cl13 …ng_cc1 -triple amdgcn -target-cpu gfx908 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX908 %s
42 // GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts…
/external/llvm-project/clang/lib/Basic/Targets/
DNVPTX.cpp201 case CudaArch::GFX908: in getTargetDefines()
/external/llvm-project/clang/test/Driver/
Damdgpu-mcpu.cl91 // RUN: %clang -### -target amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck --check-prefix=GFX908 %s
121 // GFX908: "-target-cpu" "gfx908"

12