Use StridedMemCpy in Concat/Split Kernel #4188
Conversation
| float* dest = b + b_offset * after * i; | ||
| cudaMemcpy(dest, src, len, cudaMemcpyDeviceToDevice); | ||
| } | ||
| } |
There was a problem hiding this comment.
有memory::Copy接口，可以直接用？
https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/memory/memcpy.h
| void copy_matrix<platform::CPUPlace, float>(const float* a, size_t a_offset, | ||
| float* b, size_t b_offset, | ||
| size_t len, size_t before, | ||
| size_t after) { |
There was a problem hiding this comment.
before，after分别是什么概念呢？这个函数如果要作为一个通用于Matrix的一个Copy函数，在参数命名上可能需要更Matrix化一些，比如row, col, width, height之类的。
单纯从copy_matrix的实现来看，我理解这个函数包含以下信息：
a矩阵size为X x afterb矩阵的size为Y x after- 要拷贝的子矩阵size为
before x (len / sizeof(T)) a矩阵起始位置(a_offset, 0)b矩阵起始位置(b_offset, 0)
是否类似原来的subMatrix ?
caffe2的CopyMatrix
另外，@qingqing01 我觉得这个函数不算是math操作，是否单独实现在一个比如matrix.cc的文件里比较好？matrix.cc也可以定义一些其他的通用于Matrix的函数。
There was a problem hiding this comment.
多谢 @Xreki , copy_matrix这个函数是和caffe2中的CopyMatrix类似的，把它作为一个单独的functor主要也是为了简化代码，使CPU和GPU能够公用同一个Kernel。
根据 @qingqing01 的提醒， 或许我们可以不需要这个copy_matrix，直接在Kernel里调用memory::Copy来实现内存/显存的Copy。
There was a problem hiding this comment.
根据https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/ ，可以使用cudaMemcpyAsync, 来实现异步的拷贝显存来优化性能，但这需要单独实现GPU的Kernel，这是因为cudaMemcpyAsync需要cudaStream，会造成CPU和GPU Kernel的代码不一致，不过或许可以在另外的PR来实现这个优化。
paddle/operators/split_op.h
Outdated
| const T* src = | ||
| in->data<T>() + input_offset + input_axis_dim * after * j; | ||
| memcpy(dest, src, len); | ||
| paddle::memory::Copy<Place, Place>( | ||
| boost::get<Place>(ctx.GetPlace()), static_cast<void*>(dst), |
There was a problem hiding this comment.
Maybe we can add an interface to get Place #4203
|
Please see #4205, let's give a general function for |
Yancey1989
changed the title
Yancey1989
changed the title
Fixed #4166
Fixed #3929
Fixed #3772