Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

More improvements to Conv2d speed #98

Open
wants to merge 2 commits into
base: master
Choose a base branch
from
Open

More improvements to Conv2d speed #98

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
a99d4ac
WIP: Faster Conv2d by having multiple filters per group
hunse May 8, 2016
9c491f4
WIP: flipped k, but it doesn't help
hunse May 8, 2016
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
94 changes: 66 additions & 28 deletions nengo_ocl/clra_nonlinearities.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1514,35 +1514,44 @@ def plan_conv2d(queue, X, Y, filters, biases, shape_in, shape_out,
__global ${type} *y
)
{
const int j = get_global_id(0);
const int i = get_global_id(1);
const int k = get_global_id(2);
const int j = get_global_id(1);
const int i = get_global_id(2);
const int k = get_global_id(0);
const int ij = i*${nyj} + j;

const int tj = get_local_id(0);
const int ti = get_local_id(1);
const int lsizej = get_local_size(0);
const int lsizei = get_local_size(1);
const int lsize = lsizei * lsizej;
const int tij = ti*lsizej + tj;
const int tj = get_local_id(1);
const int ti = get_local_id(2);
const int tk = get_local_id(0);
const int lsizej = get_local_size(1);
const int lsizei = get_local_size(2);
const int lsizek = get_local_size(0);

const int lsizeij = lsizei*lsizej;
const int tij = tj + ti*lsizej;

const int lsize = lsizei*lsizej*lsizek;
const int tijk = get_local_id(0) + get_local_id(1)*get_local_size(0) +
get_local_id(2)*get_local_size(0)*get_local_size(1);
const int j0 = (j - tj)*${stj} - ${pj};
const int i0 = (i - ti)*${sti} - ${pi};
__local ${type} patch[${nipatch}][${njpatch}];
% if conv:
__local ${type} filter[${si*sj}];
__local ${type} filter[${si*sj}][${nf_per}];
f += k;
% else:
f += ij;
f += ij*${nf} + k;
% endif
x += ${xstart};
y += ${ystart};

${type} out = b[k*${nyi*nyj} + ij];
const int out_ijk = k*${nyi*nyj} + i*${nyj} + j;
${type} out = b[out_ijk];

for (int c = 0; c < ${nc}; c++) {

// load image section
__global const ${type} *xc = &x[c * ${nxi * nxj}];
for (int kij = tij; kij < ${npatch}; kij += lsize) {
for (int kij = tijk; kij < ${npatch}; kij += lsize) {
const int ki = kij / ${njpatch};
const int kj = kij % ${njpatch};
const int ii = i0 + ki;
Expand All @@ -1555,27 +1564,28 @@ def plan_conv2d(queue, X, Y, filters, biases, shape_in, shape_out,

% if conv:
// load filters
__global const ${type} *fc = f + k*${nc*si*sj} + c*${si*sj};
for (int kij = tij; kij < ${si*sj}; kij += lsize) {
filter[kij] = fc[kij];
}
__global const ${type} *fc = f + c*${si*sj*nf};
for (int kij = tij; kij < ${si*sj}; kij += lsizeij)
filter[kij][tk] = fc[kij*${nf}];
% else:
__global const ${type} *filter = f + c*${si*sj*nyi*nyj*nf};
% endif
barrier(CLK_LOCAL_MEM_FENCE);

for (int ii = 0; ii < ${si}; ii++)
for (int jj = 0; jj < ${sj}; jj++)
out += patch[${sti}*ti+ii][${stj}*tj+jj]
% if conv:
out += filter[ii*${sj}+jj] * patch[${sti}*ti+ii][${stj}*tj+jj];
* filter[ii*${sj}+jj][tk];
% else:
out += f[((k*${nc} + c)*${si*sj} + ii*${sj} + jj)*${nyi*nyj}]
* patch[${sti}*ti+ii][${stj}*tj+jj];
* filter[(ii*${sj}+jj)*${nyi*nyj*nf}];
% endif

barrier(CLK_LOCAL_MEM_FENCE);
}

if (i < ${nyi} && j < ${nyj})
y[k*${nyi*nyj} + ij] = out;
if (i < ${nyi} && j < ${nyj} && k < ${nf})
y[out_ijk] = out;
}
"""

Expand All @@ -1585,12 +1595,40 @@ def plan_conv2d(queue, X, Y, filters, biases, shape_in, shape_out,
pi, pj = padding
sti, stj = strides

max_group = get_mwgs(queue, cap=256)
max_group = get_mwgs(queue, cap=128)
# max_group = get_mwgs(queue, cap=256)
assert max_group >= 32
lsize0 = min(nyj, 32)
lsize1 = min(max_group // lsize0, nyi)
lsize = (lsize0, lsize1, 1)
gsize = (round_up(nyj, lsize[0]), round_up(nyi, lsize[1]), nf)

# # lsize = (8, 8, 4)
# lsize = (4, 4, 16)
# gsize = (round_up(nyj, lsize[0]),
# round_up(nyi, lsize[1]),
# round_up(nf, lsize[2]))

lsize = (16, 4, 4)
# lsize = (32, 4, 4)
# lsize = (32, 4, 1)
gsize = (round_up(nf, lsize[0]),
round_up(nyj, lsize[1]),
round_up(nyi, lsize[2]),
)

# lsize0 = min(nyj, 32)
# lsize1 = min(max_group // lsize0, nyi, 32)
# lsize2 = min(max_group // (lsize0 * lsize1), nf)
# lsize = (lsize0, lsize1, lsize2)
# gsize = (round_up(nyj, lsize[0]),
# round_up(nyi, lsize[1]),
# round_up(nf, lsize[2]))
# print(lsize)
# print(gsize)

# lsize0 = min(nyj, 32)
# lsize1 = min(max_group // lsize0, nyi)
# lsize = (lsize0, lsize1, 1)
# gsize = (round_up(nyj, lsize[0]),
# round_up(nyi, lsize[1]),
# nf)

njpatch = (lsize[0] - 1) * stj + sj
nipatch = (lsize[1] - 1) * sti + si
Expand All @@ -1603,7 +1641,7 @@ def plan_conv2d(queue, X, Y, filters, biases, shape_in, shape_out,
textconf = dict(
type=X.ctype, conv=conv, nf=nf, nxi=nxi, nxj=nxj, nyi=nyi, nyj=nyj,
nc=nc, si=si, sj=sj, pi=pi, pj=pj, sti=sti, stj=stj,
nipatch=nipatch, njpatch=njpatch, npatch=npatch,
nipatch=nipatch, njpatch=njpatch, npatch=npatch, nf_per=lsize[2],
xstart=X.start, ystart=Y.start)
text = as_ascii(Template(text, output_encoding='ascii').render(**textconf))

Expand Down
2 changes: 1 addition & 1 deletion nengo_ocl/simulator.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -827,7 +827,7 @@ def _plan_Conv2d(self, ops):
X = self.all_data.getitem_device(self.sidx[op.input])
Y = self.all_data.getitem_device(self.sidx[op.output])
ftrans = np.asarray(np.transpose(
f, (0, 1, 2, 3) if conv else (0, 3, 4, 5, 1, 2)), order='C')
f, (1, 2, 3, 0) if conv else (3, 4, 5, 1, 2, 0)), order='C')
F = self.Array(ftrans.ravel())
B = self.Array((np.zeros(p.shape_out) + b).ravel())
plans.append(plan_conv2d(
Expand Down