haraldwolff
/
CLBlast
mirror of https://github.com/CNugteren/CLBlast.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Second version of direct reading from image tensor for convgemm: also with local memory support now

pull/319/head
Cedric Nugteren 2018-05-19 21:02:44 +02:00
parent cbcd4ff7e8
commit 27b52ac2c8
3 changed files with 122 additions and 33 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

44
src/kernels/level3/xconvgemm.opencl
View File

@ -19,35 +19,6 @@ R"(
// =================================================================================================
#if defined(ROUTINE_CONVGEMM)
// Loads global off-chip memory into thread-private register files. This function is specific for
// loading the image input tensor. This includes a bounds check.
INLINE_FUNC real GlobalToPrivateCheckedImage(const __global real* restrict imagegm, const int image_offset_batch,
const int h_id, const int w_id, const int kwg,
const int input_h, const int input_w, const int channels,
const int kernel_h, const int kernel_w,
const int pad_h, const int pad_w,
const int stride_h, const int stride_w,
const int dilation_h, const int dilation_w) {
real result;
const int kernel_2d_index = kwg % (kernel_h * kernel_w);
const int kw_id = kernel_2d_index % kernel_w;
const int kh_id = kernel_2d_index / kernel_w;
const int c_id = kwg / (kernel_h * kernel_w);
const int h_index = -pad_h + kh_id * dilation_h + stride_h * h_id;
const int w_index = -pad_w + kw_id * dilation_w + stride_w * w_id;
if (h_index >= 0 && h_index < input_h &&
w_index >= 0 && w_index < input_w) {
const int image_index = w_index + input_w * (h_index + input_h * c_id);
result = imagegm[image_index + image_offset_batch];
}
else {
SetToZero(result);
}
return result;
}
// ConvGEMM kernel
__kernel __attribute__((reqd_work_group_size(MDIMCD, NDIMCD, 1)))
void Xconvgemm(const int num_patches, const int num_kernels, const int patch_size,
@ -189,8 +160,15 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
int kwg = 0;
for (; kwg < (patch_size/WGD) * WGD; kwg+=WGD) {
// Loads data: off-chip --> local (matrix A and B)
GlobalToLocalCheckedA(colgms, alm, num_patches, col_offset_batch, kwg, false, false, num_patches, patch_size);
// Loads data: off-chip --> local
#if defined(CONVGEMM_WITH_IM2COL)
GlobalToLocalCheckedA(colgms, alm, num_patches, col_offset_batch, kwg, false, false, num_patches, patch_size);
#else
GlobalToLocalCheckedImage(imagegm, alm, image_offset_batch, h_id, w_id, kwg,
input_h, input_w, channels, kernel_h, kernel_w,
pad_h, pad_w, stride_h, stride_w,
dilation_h, dilation_w);
#endif
GlobalToLocalCheckedB(kernelgms, blm, patch_size, kernel_offset, kwg, true, false, num_kernels, patch_size);
barrier(CLK_LOCAL_MEM_FENCE);
@ -200,7 +178,7 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
for (int _pit = 0; _pit < KWID; _pit += 1) {
int kg = pwi + _pit;
// Loads data: local --> private (matrix A and B)
// Loads data: local --> private
#pragma unroll
for (int _mi = 0; _mi < MWID; _mi += 1) {
apd[_mi] = LocalToPrivateDirectA(alm, _mi, kg, false);
@ -226,7 +204,7 @@ void Xconvgemm(const int num_patches, const int num_kernels, const int patch_siz
// Loop over the remaining part (incomplete tile in K-dimension)
for (; kwg < patch_size; ++kwg) {
// Loads data: off-chip --> private (matrix A and B)
// Loads data: off-chip --> private
#pragma unroll
for (int _mi = 0; _mi < MWID; _mi += 1) {
apd[_mi] = GlobalToPrivateCheckedImage(imagegm, image_offset_batch, h_id, w_id, kwg,

110
src/kernels/levelx/xconvgemm_part1.opencl 100644
View File

@ -0,0 +1,110 @@
// =================================================================================================
// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
// width of 100 characters per line.
//
// Author(s):
// Cedric Nugteren <www.cedricnugteren.nl>
//
// This file contains the an implementation of 3D convolution on a 4D image using GEMM kernels. It
// uses parameters from the direct GEMM kernel. This is the part with the loads from memory (1/2).
//
// =================================================================================================
// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string
// literal). Comment-out this line for syntax-highlighting when developing.
R"(
// =================================================================================================
#if defined(ROUTINE_CONVGEMM)
// Loads global off-chip memory into thread-private register files. This function is specific for
// loading the image input tensor. This includes a bounds check.
INLINE_FUNC real GlobalToPrivateCheckedImage(const __global real* restrict imagegm, const int image_offset_batch,
const int h_id, const int w_id, const int kwg,
const int input_h, const int input_w, const int channels,
const int kernel_h, const int kernel_w,
const int pad_h, const int pad_w,
const int stride_h, const int stride_w,
const int dilation_h, const int dilation_w) {
// Im2col indices
const int kernel_2d_index = kwg % (kernel_h * kernel_w);
const int kw_id = kernel_2d_index % kernel_w;
const int kh_id = kernel_2d_index / kernel_w;
const int c_id = kwg / (kernel_h * kernel_w);
const int h_index = -pad_h + kh_id * dilation_h + stride_h * h_id;
const int w_index = -pad_w + kw_id * dilation_w + stride_w * w_id;
// With bounds check
real result;
if (h_index >= 0 && h_index < input_h &&
w_index >= 0 && w_index < input_w) {
const int image_index = w_index + input_w * (h_index + input_h * c_id);
result = imagegm[image_index + image_offset_batch];
}
else {
SetToZero(result);
}
return result;
}
// Loads global off-chip memory into local (shared) memory on-chip. This function is specific for
// loading the image input tensor. This includes a bounds check.
INLINE_FUNC real GlobalToLocalCheckedImage(const __global realMD* restrict imagegm, LOCAL_PTR real* alm,
const int image_offset_batch,
const int h_id, const int w_id, const int kwg,
const int input_h, const int input_w, const int channels,
const int kernel_h, const int kernel_w,
const int pad_h, const int pad_w,
const int stride_h, const int stride_w,
const int dilation_h, const int dilation_w) {
#if MDIMCD == MDIMAD
const int la0 = get_local_id(0);
const int la1 = get_local_id(1);
#else
const int tid = get_local_id(0) + MDIMCD*get_local_id(1);
const int la0 = tid % MDIMAD;
const int la1 = tid / MDIMAD;
#endif
#pragma unroll
for (int _mia = 0; _mia < MWAD; _mia += 1) {
#pragma unroll
for (int _kia = 0; _kia < KWAD; _kia += 1) {
// Computes the indices for the global memory
int mg = _mia + la0*MWAD;
int kg = _kia + la1*KWAD;
int idm = mg + GetGroupID0()*WGD;
int idk = kg + kwg;
// Im2col indices
const int kernel_2d_index = idk % (kernel_h * kernel_w);
const int kw_id = kernel_2d_index % kernel_w;
const int kh_id = kernel_2d_index / kernel_w;
const int c_id = idk / (kernel_h * kernel_w);
const int h_index = -pad_h + kh_id * dilation_h + stride_h * h_id;
const int w_index = -pad_w + kw_id * dilation_w + stride_w * w_id;
// Loads the data from global memory into the local memory
if (h_index >= 0 && h_index < input_h &&
w_index >= 0 && w_index < input_w) {
const int image_index = w_index + input_w * (h_index + input_h * c_id);
const real result = imagegm[image_index + image_offset_batch];
alm[kg*(WGD + PADA) + mg] = result;
}
else {
SetToZero(alm[kg*(WGD + PADA) + mg]);
}
}
}
}
#endif
// =================================================================================================
// End of the C++11 raw string literal
)"
// =================================================================================================

1
src/routines/levelx/xconvgemm.cpp
View File

@ -31,6 +31,7 @@ Xconvgemm<T>::Xconvgemm(Queue &queue, EventPointer event, const std::string &nam
#include "../../kernels/level3/xgemm_direct_part2.opencl"
#include "../../kernels/level3/xgemm_direct_part3.opencl"
, // separated in multiple parts to prevent C1091 in MSVC 2013
#include "../../kernels/levelx/xconvgemm_part1.opencl"
#include "../../kernels/level3/xconvgemm.opencl"
}) {
}