提交 03ccb9a4 编写于 作者: Y Yihua Xu

Optimize the stack operator

上级 7a64d48f
...@@ -147,16 +147,23 @@ class StackKernel : public framework::OpKernel<T> { ...@@ -147,16 +147,23 @@ class StackKernel : public framework::OpKernel<T> {
auto &dim = x[0]->dims(); auto &dim = x[0]->dims();
for (auto i = 0; i < axis; ++i) pre *= dim[i]; for (auto i = 0; i < axis; ++i) pre *= dim[i];
for (auto i = axis; i < dim.size(); ++i) post *= dim[i]; for (auto i = axis; i < dim.size(); ++i) post *= dim[i];
int total_num = pre * n * post;
auto &dev_ctx = ctx.template device_context<DeviceContext>();
#ifdef __NVCC__ #ifdef __NVCC__
thrust::device_vector<const T *> device_x_vec(x_datas); thrust::device_vector<const T *> device_x_vec(x_datas);
auto x_data_arr = device_x_vec.data().get(); auto x_data_arr = device_x_vec.data().get();
#else #else
auto x_data_arr = x_datas.data(); auto x_data_arr = x_datas.data();
#endif #endif
StackFunctorForRange(dev_ctx, x_data_arr, y_data, total_num, n, post); size_t x_offset = 0;
size_t y_offset = 0;
for (int i = 0; i < pre; i++) {
for (int j = 0; j < n; j++) {
std::memcpy(y_data + y_offset, x_data_arr[j] + x_offset,
post * sizeof(T));
y_offset += post;
}
x_offset += post;
}
#ifdef __NVCC__ #ifdef __NVCC__
// Wait() must be called because device_x_vec may be destructed before // Wait() must be called because device_x_vec may be destructed before
// kernel ends // kernel ends
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册