Home
last modified time | relevance | path

Searched refs:CODE (Results 1 – 25 of 356) sorted by relevance

12345678910>>...15

/external/llvm-project/polly/test/GPGPU/
Dscalar-parameter.ll3 ; RUN: FileCheck -check-prefix=CODE %s
19 ; CODE: Code
20 ; CODE-NEXT: ====
21 ; CODE-NEXT: # host
22 ; CODE-NEXT: {
23 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(float), cudaMemcp…
24 ; CODE-NEXT: {
25 ; CODE-NEXT: dim3 k0_dimBlock(32);
26 ; CODE-NEXT: dim3 k0_dimGrid(32);
27 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, MemRef_b);
[all …]
Dnon-read-only-scalars.ll3 ; RUN: FileCheck -check-prefix=CODE %s
34 ; CODE: dim3 k0_dimBlock(32);
35 ; CODE-NEXT: dim3 k0_dimGrid(1);
36 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
37 ; CODE-NEXT: cudaCheckKernel();
38 ; CODE-NEXT: }
40 ; CODE: {
41 ; CODE-NEXT: dim3 k1_dimBlock;
42 ; CODE-NEXT: dim3 k1_dimGrid;
43 ; CODE-NEXT: kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_sum_0__phi);
[all …]
Dmostly-sequential.ll3 ; RUN: FileCheck -check-prefix=CODE %s
16 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (128) * sizeof(float), cudaMemcpy…
17 ; CODE-NEXT: {
18 ; CODE-NEXT: dim3 k0_dimBlock(32);
19 ; CODE-NEXT: dim3 k0_dimGrid(4);
20 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
21 ; CODE-NEXT: cudaCheckKernel();
22 ; CODE-NEXT: }
24 ; CODE: {
25 ; CODE-NEXT: dim3 k1_dimBlock;
[all …]
Dparametric-loop-bound.ll3 ; RUN: FileCheck -check-prefix=CODE %s
16 ; CODE: if (n >= 1) {
17 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (n) * sizeof(i64), cudaMemcpyHost…
18 ; CODE-NEXT: {
19 ; CODE-NEXT: dim3 k0_dimBlock(32);
20 ; CODE-NEXT: dim3 k0_dimGrid(n >= 1048545 ? 32768 : (n + 31) / 32);
21 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, n);
22 ; CODE-NEXT: cudaCheckKernel();
23 ; CODE-NEXT: }
25 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (n) * sizeof(i64), cudaMemcpyDevi…
[all …]
Dhost-statement.ll4 ; RUN: FileCheck -check-prefix=CODE %s
21 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (512) * (512) * sizeof(double), c…
22 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_R, MemRef_R, (p_0 + 1) * (512) * sizeof(double…
23 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_Q, MemRef_Q, (512) * (512) * sizeof(double), c…
24 ; CODE-NEXT: {
25 ; CODE-NEXT: dim3 k0_dimBlock(32);
26 ; CODE-NEXT: dim3 k0_dimGrid(16);
27 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p…
28 ; CODE-NEXT: cudaCheckKernel();
29 ; CODE-NEXT: }
[all …]
Dsize-cast.ll3 ; RUN: FileCheck -check-prefix=CODE %s
12 ; CODE: if (arg >= 1 && arg1 == 0) {
13 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_arg3, MemRef_arg3, (arg) * sizeof(double), cud…
14 ; CODE-NEXT: {
15 ; CODE-NEXT: dim3 k0_dimBlock(32);
16 ; CODE-NEXT: dim3 k0_dimGrid(arg >= 1048545 ? 32768 : (arg + 31) / 32);
17 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_arg3, dev_MemRef_arg2, arg, arg1…
18 ; CODE-NEXT: cudaCheckKernel();
19 ; CODE-NEXT: }
21 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_arg2, dev_MemRef_arg2, (arg) * sizeof(double), cudaMemc…
[all …]
Dnon-zero-array-offset.ll3 ; RUN: FileCheck -check-prefix=CODE %s
10 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (16) * sizeof(float), cudaMemcpyHos…
11 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (8) * sizeof(float), cudaMemcpyHost…
13 ; CODE: dim3 k0_dimBlock(8);
14 ; CODE-NEXT: dim3 k0_dimGrid(1);
15 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
16 ; CODE-NEXT: cudaCheckKernel();
17 ; CODE-NEXT: }
19 ; CODE: {
20 ; CODE-NEXT: dim3 k1_dimBlock(8);
[all …]
Dscheduler-timeout.ll3 ; RUN: FileCheck -check-prefix=CODE %s
30 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (4096) * (4096) * sizeof(float), …
31 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (4096) * (4096) * sizeof(float), …
32 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_D, MemRef_D, (4096) * (4096) * sizeof(float), …
33 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_C, MemRef_C, (4096) * (4096) * sizeof(float), …
34 ; CODE-NEXT: {
35 ; CODE-NEXT: dim3 k0_dimBlock(16, 32);
36 ; CODE-NEXT: dim3 k0_dimGrid(128, 128);
37 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_tmp, dev_MemRef_A, MemRef_alpha,…
38 ; CODE-NEXT: cudaCheckKernel();
[all …]
Dmanaged-pointers-preparation.ll8 ; RUN: -polly-acc-dump-code < %s | FileCheck %s -check-prefix=CODE
20 ; CODE: if (p_0_loaded_from___data_runcontrol_MOD_lmulti_layer == 0) {
21 ; CODE-NEXT: {
22 ; CODE-NEXT: dim3 k0_dimBlock;
23 ; CODE-NEXT: dim3 k0_dimGrid;
24 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef__pn__phi, p_0_loaded_from___dat…
25 ; CODE-NEXT: cudaCheckKernel();
26 ; CODE-NEXT: }
28 ; CODE: } else {
29 ; CODE-NEXT: {
[all …]
Dinvalid-kernel.ll3 ; RUN: FileCheck -check-prefix=CODE %s
23 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (1024) * sizeof(i64), cudaMemcpyH…
24 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i64), cudaMemcpyH…
25 ; CODE-NEXT: {
26 ; CODE-NEXT: dim3 k0_dimBlock(32);
27 ; CODE-NEXT: dim3 k0_dimGrid(32);
28 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_B, dev_MemRef_A);
29 ; CODE-NEXT: cudaCheckKernel();
30 ; CODE-NEXT: }
32 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i64), cudaMemcpyDevice…
[all …]
Dregion-stmt.ll3 ; RUN: FileCheck -check-prefix=CODE %s
8 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (128) * sizeof(float), cudaMemcpy…
9 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (128) * sizeof(float), cudaMemcpy…
10 ; CODE-NEXT: {
11 ; CODE-NEXT: dim3 k0_dimBlock(32);
12 ; CODE-NEXT: dim3 k0_dimGrid(4);
13 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, dev_MemRef_B);
14 ; CODE-NEXT: cudaCheckKernel();
15 ; CODE-NEXT: }
17 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_B, dev_MemRef_B, (128) * sizeof(float), cudaMemcpyDevic…
[all …]
Dshared-memory-two-dimensional.ll4 ; RUN: FileCheck -check-prefix=CODE %s
21 ; CODE: # kernel0
22 ; CODE-NEXT: {
23 ; CODE-NEXT: if (t0 <= 7)
24 ; CODE-NEXT: for (int c0 = 0; c0 <= 15; c0 += 1)
25 ; CODE-NEXT: read(c0, t0);
26 ; CODE-NEXT: read(t0);
27 ; CODE-NEXT: sync0();
28 ; CODE-NEXT: for (int c3 = 0; c3 <= 15; c3 += 1)
29 ; CODE-NEXT: for (int c4 = 0; c4 <= 7; c4 += 1)
[all …]
Dphi-nodes-in-kernel.ll3 ; RUN: FileCheck -check-prefix=CODE %s
35 ; CODE: cudaCheckReturn(cudaMalloc((void **) &dev_MemRef_c, (50) * sizeof(i32)));
37 ; CODE: {
38 ; CODE-NEXT: dim3 k0_dimBlock(32);
39 ; CODE-NEXT: dim3 k0_dimGrid(2);
40 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_c);
41 ; CODE-NEXT: cudaCheckKernel();
42 ; CODE-NEXT: }
44 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_c, dev_MemRef_c, (50) * sizeof(i32), cudaMemcpyDevi…
45 ; CODE-NEXT: cudaCheckReturn(cudaFree(dev_MemRef_c));
[all …]
/external/llvm-project/llvm/test/Analysis/CostModel/AArch64/
Dmin-max.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
7 ; CODE-LABEL: umin.v8i8
8 ; CODE: bb.0
9 ; CODE-NEXT: umin v{{.*}}.8b, v{{.*}}.8b, v{{.*}}.8b
10 ; CODE-NEXT: ret
21 ; CODE-LABEL: umin.v9i8
22 ; CODE: bb.0
23 ; CODE-NEXT: umin v{{.*}}.16b, v{{.*}}.16b, v{{.*}}.16b
24 ; CODE-NEXT: ret
35 ; CODE-LABEL: umin.v4i16
[all …]
Dvector-select.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
8 ; CODE-LABEL: v8i8_select_eq
9 ; CODE: bb.0
10 ; CODE-NEXT: cmeq v{{.+}}.8b, v{{.+}}.8b, v{{.+}}.8b
11 ; CODE-NEXT: bif v{{.+}}.8b, v{{.+}}.8b, v{{.+}}.8b
12 ; CODE-NEXT: ret
24 ; CODE-LABEL: v16i8_select_sgt
25 ; CODE: bb.0
26 ; CODE-NEXT: cmgt v{{.+}}.16b, v{{.+}}.16b, v{{.+}}.16b
27 ; CODE-NEXT: bif v{{.+}}.16b, v{{.+}}.16b, v{{.+}}.16b
[all …]
Dvector-reduce.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
6 ; CODE-LABEL: add.i8.v8i8
7 ; CODE: addv b0, v0.8b
15 ; CODE-LABEL: add.i8.v16i8
16 ; CODE: addv b0, v0.16b
24 ; CODE-LABEL: add.i16.v4i16
25 ; CODE: addv h0, v0.4h
33 ; CODE-LABEL: add.i16.v8i16
34 ; CODE: addv h0, v0.8h
42 ; CODE-LABEL: add.i32.v4i32
[all …]
Dfree-widening-casts.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
7 ; CODE-LABEL: uaddl_8h
8 ; CODE: uaddl v0.8h, v0.8b, v1.8b
19 ; CODE-LABEL: uaddl_4s
20 ; CODE: uaddl v0.4s, v0.4h, v1.4h
31 ; CODE-LABEL: uaddl_2d
32 ; CODE: uaddl v0.2d, v0.2s, v1.2s
43 ; CODE-LABEL: uaddl2_8h
44 ; CODE: uaddl2 v2.8h, v0.16b, v1.16b
45 ; CODE-NEXT: uaddl v0.8h, v0.8b, v1.8b
[all …]
Dshuffle-transpose.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
6 ; CODE-LABEL: trn1.v8i8
7 ; CODE: trn1 v0.8b, v0.8b, v1.8b
15 ; CODE-LABEL: trn2.v8i8
16 ; CODE: trn2 v0.8b, v0.8b, v1.8b
24 ; CODE-LABEL: trn1.v16i8
25 ; CODE: trn1 v0.16b, v0.16b, v1.16b
33 ; CODE-LABEL: trn2.v16i8
34 ; CODE: trn2 v0.16b, v0.16b, v1.16b
42 ; CODE-LABEL: trn1.v4i16
[all …]
Dshuffle-select.ll2 ; RUN: llc < %s -mtriple=aarch64--linux-gnu | FileCheck %s --check-prefix=CODE
6 ; CODE-LABEL: sel.v8i8
7 ; CODE: tbl v0.8b, { v0.16b }, v2.8b
15 ; CODE-LABEL: sel.v16i8
16 ; CODE: tbl v0.16b, { v0.16b, v1.16b }, v2.16b
24 ; CODE-LABEL: sel.v4i16
25 ; CODE: rev32 v0.4h, v0.4h
26 ; CODE: trn2 v0.4h, v0.4h, v1.4h
34 ; CODE-LABEL: sel.v8i16
35 ; CODE: tbl v0.16b, { v0.16b, v1.16b }, v2.16b
[all …]
/external/llvm-project/llvm/test/CodeGen/ARM/Windows/
Dvla.ll2 ; RUN: | FileCheck %s -check-prefix CHECK-SMALL-CODE
4 ; RUN: | FileCheck %s -check-prefix CHECK-LARGE-CODE
6 ; RUN: | FileCheck %s -check-prefix CHECK-SMALL-CODE
16 ; CHECK-SMALL-CODE: adds [[R4:r[0-9]+]], #7
17 ; CHECK-SMALL-CODE: bic [[R4]], [[R4]], #4
18 ; CHECK-SMALL-CODE: lsrs r4, [[R4]], #2
19 ; CHECK-SMALL-CODE: bl __chkstk
20 ; CHECK-SMALL-CODE: sub.w sp, sp, r4
22 ; CHECK-LARGE-CODE: adds [[R4:r[0-9]+]], #7
23 ; CHECK-LARGE-CODE: bic [[R4]], [[R4]], #4
[all …]
Dchkstk.ll2 ; RUN: | FileCheck -check-prefix CHECK-DEFAULT-CODE-MODEL %s
5 ; RUN: | FileCheck -check-prefix CHECK-LARGE-CODE-MODEL %s
13 ; CHECK-DEFAULT-CODE-MODEL: check_watermark:
14 ; CHECK-DEFAULT-CODE-MODEL: movw r4, #1024
15 ; CHECK-DEFAULT-CODE-MODEL: bl __chkstk
16 ; CHECK-DEFAULT-CODE-MODEL: sub.w sp, sp, r4
18 ; CHECK-LARGE-CODE-MODEL: check_watermark:
19 ; CHECK-LARGE-CODE-MODEL-DAG: movw r12, :lower16:__chkstk
20 ; CHECK-LARGE-CODE-MODEL-DAG: movt r12, :upper16:__chkstk
21 ; CHECK-LARGE-CODE-MODEL-DAG: movw r4, #1024
[all …]
Dstack-probe-non-default.ll2 ; RUN: | FileCheck %s -check-prefix CHECK-DEFAULT-CODE-MODEL
5 ; RUN: | FileCheck %s -check-prefix CHECK-LARGE-CODE-MODEL
21 ; CHECK-DEFAULT-CODE-MODEL-NOT: __chkstk
22 ; CHECK-DEFAULT-CODE-MODEL: sub.w sp, sp, #4096
24 ; CHECK-LARGE-CODE-MODEL-NOT: movw r12, :lower16:__chkstk
25 ; CHECK-LARGE-CODE-MODEL-NOT: movt r12, :upper16:__chkstk
26 ; CHECK-LARGE-CODE-MODEL: sub.w sp, sp, #4096
/external/llvm/test/CodeGen/ARM/Windows/
Dvla.ll2 ; RUN: | FileCheck %s -check-prefix CHECK-SMALL-CODE
4 ; RUN: | FileCheck %s -check-prefix CHECK-LARGE-CODE
16 ; CHECK-SMALL-CODE: adds [[R4:r[0-9]+]], #7
17 ; CHECK-SMALL-CODE: bic [[R4]], [[R4]], #7
18 ; CHECK-SMALL-CODE: lsrs r4, [[R4]], #2
19 ; CHECK-SMALL-CODE: bl __chkstk
20 ; CHECK-SMALL-CODE: sub.w sp, sp, r4
22 ; CHECK-LARGE-CODE: adds [[R4:r[0-9]+]], #7
23 ; CHECK-LARGE-CODE: bic [[R4]], [[R4]], #7
24 ; CHECK-LARGE-CODE: lsrs r4, [[R4]], #2
[all …]
Dchkstk.ll2 ; RUN: | FileCheck -check-prefix CHECK-DEFAULT-CODE-MODEL %s
5 ; RUN: | FileCheck -check-prefix CHECK-LARGE-CODE-MODEL %s
13 ; CHECK-DEFAULT-CODE-MODEL: check_watermark:
14 ; CHECK-DEFAULT-CODE-MODEL: movw r4, #1024
15 ; CHECK-DEFAULT-CODE-MODEL: bl __chkstk
16 ; CHECK-DEFAULT-CODE-MODEL: sub.w sp, sp, r4
18 ; CHECK-LARGE-CODE-MODEL: check_watermark:
19 ; CHECK-LARGE-CODE-MODEL: movw r12, :lower16:__chkstk
20 ; CHECK-LARGE-CODE-MODEL: movt r12, :upper16:__chkstk
21 ; CHECK-LARGE-CODE-MODEL: movw r4, #1024
[all …]
/external/llvm-project/llvm/test/CodeGen/AArch64/
Dchkstk.ll2 ; RUN: | FileCheck -check-prefix CHECK-DEFAULT-CODE-MODEL %s
7 ; RUN: | FileCheck -check-prefix CHECK-LARGE-CODE-MODEL %s
17 ; CHECK-DEFAULT-CODE-MODEL: check_watermark:
18 ; CHECK-DEFAULT-CODE-MODEL-DAG: stp x29, x30, [sp
19 ; CHECK-DEFAULT-CODE-MODEL-DAG: mov x15, #256
20 ; CHECK-DEFAULT-CODE-MODEL: bl __chkstk
21 ; CHECK-DEFAULT-CODE-MODEL: sub sp, sp, x15, lsl #4
25 ; CHECK-LARGE-CODE-MODEL: check_watermark:
26 ; CHECK-LARGE-CODE-MODEL-DAG: stp x29, x30, [sp
27 ; CHECK-LARGE-CODE-MODEL-DAG: mov x15, #256
[all …]

12345678910>>...15