其他
基于OneFlow实现Unfold、Fold算子
撰文 | zzk
1
从卷积层说起
from mxnet.gluon import nn
def corr2d(X, K): # 本函数已保存在d2lzh包中方便以后使用
h, w = K.shape
Y = nd.zeros((X.shape[0] - h + 1, X.shape[1] - w + 1))
for i in range(Y.shape[0]):
for j in range(Y.shape[1]):
Y[i, j] = (X[i: i + h, j: j + w] * K).sum()
return Y
2
初见img2col
https://github.com/microsoft/AI-System/blob/main/docs/SystemforAI-4-Computer%20architecture%20for%20Matrix%20computation.pdf 这是微软的AISystem仓库对应的章节,强烈推荐大家去学习(本人鸽了好久没看完)
3
理解img2col
假设输入特征图为(N, Cin, H, W),卷积核参数为(K, K, Cin, Cout), 输出特征图的长宽为(Oh, Ow)
(N, Cin*K*K, Oh*Ow)
这么一个三维向量。(Cout, K*K*Cin)
。(N, Cout, Oh*Ow)
,这也是我们预期的卷积结果。4
img2col源码
out_h = int((height + 2*pad - ksize) / stride) + 1
out_w = int((width + 2*pad - ksize) / stride) + 1
channels_cols = channels*ksize*ksize
out_shape = (channels_cols, out_h*out_w)
elem_cnt = out_shape[0] * out_shape[1]
out_array = np.zeros(shape=elem_cnt, dtype=np.float32)
out_h
和out_w
就是输出的高,宽,采用的是卷积输出的公式:channel_cols
则是之前我们提到的,img2col会把第二个维度变换为C_in*K*K
。channel_cols
的for循环# 分别计算一个k*k的窗口内,h,w的偏移量
kh_offset = (c // ksize) % ksize
kw_offset = c % ksize
# 计算当前处理的通道index
c_im = c // ksize // ksize
for w in range(out_w):
im_row = kh_offset + h * stride
im_col = kw_offset + w * stride
index = (c * out_h + h) * out_w + w
out_array = np.reshape(out_array, out_shape)
return out_array
img2col_get_pixel
是一个合法取元素的函数,如果出现越界范围(比如小于0,或者大于Oh),那么就是取到padding的部分,此时我们返回0。row -= pad
col -= pad
if row < 0 or col < 0 or row >= height or col >= width:
return 0
return im[int(col + width * (row + height * channel))] # batch*w*h*c + width*height*channel + width*row + col
out = darknet_img2col(x, channels=1, height=3, width=3, ksize=2, stride=1, pad=0)
[2. 3. 5. 6.]
[4. 5. 7. 8.]
[5. 6. 8. 9.]]
col2img
OneFlow对应的实现
5
OneFlow版本的Unfold
NdIndexHelper
类,构造的时候我们可以传入高维shape,然后调用OffsetToNdIndex
来进行一维offset到高维索引的转换。NdIndexHelper
out_index_helper = NdIndexOffsetHelper<INDEX_T, kOutputNDim>(output_dims); // 格式为(N, C, Kh, Kw, Oh, Ow)
// index_b format: (N, C, D, H, W) or (N, D, H, W, C)
// return: true indicates out-of-bound, otherwise in-bound
template<typename INDEX_T, int NDIM, int SDIM>
OF_DEVICE_FUNC bool UnfoldIndexTransform(const UnfoldParams<INDEX_T, NDIM, SDIM>& params,
const INDEX_T* index_a, INDEX_T* index_b) {
// batch dim index transform
index_b[0] = index_a[0];
// channel dim index transform
using ParamType = UnfoldParams<INDEX_T, NDIM, SDIM>;
index_b[ParamType::kInputChannelDim] = index_a[ParamType::kOutputChannelDim];
// spatial dim index transform
// D,H,W spatial dim index transform
for (int64_t d = 0; d < NDIM; ++d) {
INDEX_T idx = index_a[SDIM + NDIM + d] * params.stride[d]
+ index_a[SDIM + d] * params.dilation[d] - params.padding[d];
if (idx < 0 || idx >= params.dims[d]) return true;
index_b[SDIM + d] = idx;
}
return false;
}
模板参数 INDEX_T表示Index的数据类型(可以有int32_t, int64_t),NDIM表示处理几维数据(这里我们是2维),SDIM则是决定通道维所在位置,SDIM=1是NHWC格式,SDIM=2则是NCHW格式(这里我们取2) 输入参数 index_a表示输出的NdIndexHelper,index_b则表示的是输入的NdIndexHelper
index_b
即输入的NdIndexHelper,且return false。__global__ void CudaUnfoldForward(UnfoldParams<INDEX_T, NDIM, SDIM> params, const T* in, T* out) {
CUDA_1D_KERNEL_LOOP_T(INDEX_T, out_offset, params.out_elem_cnt) {
using ParamType = UnfoldParams<INDEX_T, NDIM, SDIM>;
INDEX_T in_index[ParamType::kInputNDim] = {0};
INDEX_T out_index[ParamType::kOutputNDim] = {0};
params.out_index_helper.OffsetToNdIndex(out_offset, out_index);
if (!UnfoldIndexTransform<INDEX_T, NDIM, SDIM>(params, out_index, in_index)) {
INDEX_T in_offset = params.in_index_helper.NdIndexToOffset(in_index);
out[out_offset] = in[in_offset];
} else {
out[out_offset] = static_cast<T>(kUnfoldPaddingValue);
}
}
}
首先根据offset来计算当前处理输出元素的NdIndex 然后判断UnfoldIndexTransform该方法的返回值 如果为false,则说明我们可以取到输入元素,将其index转换为1d的offset,并赋值给输出 如果为true,则越界,我们填充先前设定好的一个padding_value(0)
6
OneFlow版本的Fold
__global__ void CudaFoldForward(FoldParams<INDEX_T, NDIM, SDIM> params, const T* input_ptr,
T* output_ptr) {
CUDA_1D_KERNEL_LOOP_T(INDEX_T, in_offset, params.in_elem_cnt) {
using ParamType = FoldParams<INDEX_T, NDIM, SDIM>;
INDEX_T in_index[ParamType::kInputNDim] = {0};
INDEX_T out_index[ParamType::kOutputNDim] = {0};
params.in_index_helper.OffsetToNdIndex(in_offset, in_index);
if (!FoldIndexTransform<INDEX_T, NDIM, SDIM>(params, in_index, out_index)) {
INDEX_T out_offset = params.out_index_helper.NdIndexToOffset(out_index);
XPUAdd<T>::Invoke(&input_ptr[in_offset], &output_ptr[out_offset]);
} else {
continue;
}
}
}
FoldIndexTransform
返回的是false,则计算输出的offset,并使用原子加atomic add,把输入元素累加到该输出位置。