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
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
|
/*
* Subgroup scan inclusive add functions
*/
#ifndef HALF
kernel void compiler_subgroup_scan_inclusive_add_short(global short *src, global short *dst) {
short val = src[get_global_id(0)];
short sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_ushort(global ushort *src, global ushort *dst) {
ushort val = src[get_global_id(0)];
ushort sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_uint(global uint *src, global uint *dst) {
uint val = src[get_global_id(0)];
uint sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_long(global long *src, global long *dst) {
long val = src[get_global_id(0)];
long sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_ulong(global ulong *src, global ulong *dst) {
ulong val = src[get_global_id(0)];
ulong sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_add_float(global float *src, global float *dst) {
float val = src[get_global_id(0)];
float sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
/*
* Subgroup scan inclusive max functions
*/
kernel void compiler_subgroup_scan_inclusive_max_short(global short *src, global short *dst) {
short val = src[get_global_id(0)];
short sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_ushort(global ushort *src, global ushort *dst) {
ushort val = src[get_global_id(0)];
ushort sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_uint(global uint *src, global uint *dst) {
uint val = src[get_global_id(0)];
uint sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_long(global long *src, global long *dst) {
long val = src[get_global_id(0)];
long sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_ulong(global ulong *src, global ulong *dst) {
ulong val = src[get_global_id(0)];
ulong sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_float(global float *src, global float *dst) {
float val = src[get_global_id(0)];
float sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
/*
* Subgroup scan inclusive min functions
*/
kernel void compiler_subgroup_scan_inclusive_min_short(global short *src, global short *dst) {
short val = src[get_global_id(0)];
short sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_ushort(global ushort *src, global ushort *dst) {
ushort val = src[get_global_id(0)];
ushort sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_int(global int *src, global int *dst) {
int val = src[get_global_id(0)];
int sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_uint(global uint *src, global uint *dst) {
uint val = src[get_global_id(0)];
uint sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_long(global long *src, global long *dst) {
long val = src[get_global_id(0)];
long sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_ulong(global ulong *src, global ulong *dst) {
ulong val = src[get_global_id(0)];
ulong sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_float(global float *src, global float *dst) {
float val = src[get_global_id(0)];
float sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
#else
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
kernel void compiler_subgroup_scan_inclusive_add_half(global half *src, global half *dst) {
half val = src[get_global_id(0)];
half sum = sub_group_scan_inclusive_add(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_max_half(global half *src, global half *dst) {
half val = src[get_global_id(0)];
half sum = sub_group_scan_inclusive_max(val);
dst[get_global_id(0)] = sum;
}
kernel void compiler_subgroup_scan_inclusive_min_half(global half *src, global half *dst) {
half val = src[get_global_id(0)];
half sum = sub_group_scan_inclusive_min(val);
dst[get_global_id(0)] = sum;
}
#endif
|
