1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
|
;=========================== begin_copyright_notice ============================
;
; Copyright (C) 2023-2024 Intel Corporation
;
; SPDX-License-Identifier: MIT
;
;============================ end_copyright_notice =============================
; REQUIRES: llvm-14-plus
; RUN: igc_opt --opaque-pointers %s -S -o - -igc-dpas-scan -igc-serialize-metadata -platformpvc | FileCheck %s
; CHECK: !{!"DisableEUFusion", i1 true}
; .cl source
; float __builtin_IB_sub_group_fdpas_hf_hf_8_1 (float acc, int a, int8 b) __attribute__((const));
;__kernel void SimpleArg(float acc, int src1, int8 src2, __global float *dst) {
; for (int i = 0; i < 3; i++) {
; dst[i] = __builtin_IB_sub_group_fdpas_hf_hf_8_1(acc, src1, src2);
; if (dst[i] < 0.f) break;
; }
;}
declare spir_func float @__builtin_IB_sub_group_fdpas_hf_hf_8_1(float, i32, <8 x i32>) #0
define spir_kernel void @SimpleArg(float %acc, i32 %src1, <8 x i32> %src2, float addrspace(1)* %dst) #1 {
entry:
%dpas = call float @llvm.genx.GenISA.sub.group.dpas.f32.f32.i32.v8i32(float %acc, i32 %src1, <8 x i32> %src2, i32 12, i32 12, i32 8, i32 1, i1 false)
store float %dpas, float addrspace(1)* %dst, align 4
%cmp3 = fcmp olt float %dpas, 0.000000e+00
br i1 %cmp3, label %entry.for.end_crit_edge, label %for.cond
for.cond: ; preds = %entry
%dpas.1 = call float @llvm.genx.GenISA.sub.group.dpas.f32.f32.i32.v8i32(float %acc, i32 %src1, <8 x i32> %src2, i32 12, i32 12, i32 8, i32 1, i1 false)
%0 = inttoptr i32 4 to float addrspace(131072)*
store float %dpas.1, float addrspace(131072)* %0, align 4
%cmp3.1 = fcmp olt float %dpas.1, 0.000000e+00
br i1 %cmp3.1, label %for.cond.for.end_crit_edge, label %for.cond.1
for.cond.1: ; preds = %for.cond
%dpas.2 = call float @llvm.genx.GenISA.sub.group.dpas.f32.f32.i32.v8i32(float %acc, i32 %src1, <8 x i32> %src2, i32 12, i32 12, i32 8, i32 1, i1 false)
%1 = inttoptr i32 8 to float addrspace(131072)*
store float %dpas.2, float addrspace(131072)* %1, align 8
br label %UnifiedReturnBlock
for.cond.for.end_crit_edge: ; preds = %for.cond
br label %for.end
entry.for.end_crit_edge: ; preds = %entry
br label %for.end
for.end: ; preds = %for.cond.for.end_crit_edge, %entry.for.end_crit_edge
br label %UnifiedReturnBlock
UnifiedReturnBlock: ; preds = %for.cond.1, %for.end
ret void
}
declare float @llvm.genx.GenISA.sub.group.dpas.f32.f32.i32.v8i32(float, i32, <8 x i32>, i32, i32, i32, i32, i1) #1
!68 = !{!"DisableEUFusion", i1 false}
!igc.functions = !{!332}
!332 = !{void (float, i32, <8 x i32>, float addrspace(1)*)* @SimpleArg, !333}
!333 = !{!334}
!334 = !{!"function_type", i32 0}
|
