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
|
#ifndef __OPENCL_VERSION__
#include "header.cl"
#endif
__kernel void TEMPLATE(bias_forward,Dtype)(const int_tp n,
__global const Dtype* in,
__global const Dtype* bias,
const int_tp bias_dim,
const int_tp inner_dim,
__global Dtype* out) {
for (int_tp index = get_global_id(0); index < n;
index += get_global_size(0)) {
const int_tp bias_index = (index / inner_dim) % bias_dim;
out[index] = in[index] + bias[bias_index];
}
}
__kernel void TEMPLATE(scale_forward,Dtype)(const int_tp n,
__global const Dtype* in,
__global const Dtype* scale,
const int_tp scale_dim,
const int_tp inner_dim,
__global Dtype* out) {
for (int_tp index = get_global_id(0); index < n;
index += get_global_size(0)) {
const int_tp scale_index = (index / inner_dim) % scale_dim;
out[index] = in[index] * scale[scale_index];
}
}
__kernel void TEMPLATE(scale_bias_forward,Dtype)(const int_tp n,
__global const Dtype* in,
__global const Dtype* scale,
__global const Dtype* bias,
const int_tp scale_dim,
const int_tp inner_dim,
__global Dtype* out) {
for (int_tp index = get_global_id(0); index < n;
index += get_global_size(0)) {
const int_tp scale_index = (index / inner_dim) % scale_dim;
out[index] = in[index] * scale[scale_index] + bias[scale_index];
}
}
|
