Liu Yiqun
8 years ago
commit
8495f3f04a
Binary file not shown.
|
Before Width: | Height: | Size: 24 KiB After Width: | Height: | Size: 24 KiB |
@ -0,0 +1,52 @@
|
||||
/*
|
||||
Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License.
|
||||
*/
|
||||
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include "paddle/framework/lod_tensor.h"
|
||||
#include "paddle/platform/assert.h"
|
||||
|
||||
#include <gtest/gtest.h>
|
||||
|
||||
__global__ void test(size_t* a, int size) {
|
||||
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size;
|
||||
i += blockDim.x * gridDim.x) {
|
||||
a[i] *= 2;
|
||||
}
|
||||
}
|
||||
|
||||
TEST(LoDTensor, LoDInGPU) {
|
||||
paddle::framework::Tensor tensor;
|
||||
paddle::framework::LoDTensor lod_tensor;
|
||||
paddle::platform::GPUPlace place(0);
|
||||
|
||||
paddle::framework::LoD src_lod;
|
||||
src_lod.push_back(std::vector<size_t>{0, 2, 4, 6, 8, 10, 12, 14});
|
||||
|
||||
tensor.Resize({14, 16});
|
||||
tensor.mutable_data<float>(place);
|
||||
|
||||
lod_tensor.set_lod(src_lod);
|
||||
lod_tensor.set_tensor(&tensor);
|
||||
CHECK_EQ(lod_tensor.lod_element(0, 2), 4);
|
||||
CHECK_EQ(lod_tensor.lod_element(0, 4), 8);
|
||||
|
||||
auto lod = lod_tensor.lod();
|
||||
|
||||
test<<<1, 8>>>(lod[0].data(), lod[0].size());
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
for (size_t i = 0; i < src_lod[0].size(); ++i) {
|
||||
CHECK_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2);
|
||||
}
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,79 @@
|
||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License. */
|
||||
|
||||
#include "paddle/operators/concat_op.h"
|
||||
#include <vector>
|
||||
|
||||
namespace paddle {
|
||||
namespace operators {
|
||||
using framework::Tensor;
|
||||
|
||||
class ConcatOp : public framework::OperatorWithKernel {
|
||||
public:
|
||||
using framework::OperatorWithKernel::OperatorWithKernel;
|
||||
|
||||
protected:
|
||||
void InferShape(const framework::InferShapeContext &ctx) const override {
|
||||
auto ins = ctx.MultiInput<framework::Tensor>("X");
|
||||
auto *out = ctx.Output<framework::Tensor>("Out");
|
||||
size_t axis = static_cast<size_t>(ctx.Attr<int>("axis"));
|
||||
size_t n = ins.size();
|
||||
|
||||
PADDLE_ENFORCE_GT(n, 1, "Input tensors count should > 1.");
|
||||
|
||||
auto out_dims = ins[0]->dims();
|
||||
size_t in_zero_dims_size = out_dims.size();
|
||||
for (size_t i = 1; i < n; i++) {
|
||||
for (size_t j = 0; j < in_zero_dims_size; j++) {
|
||||
if (j == axis) {
|
||||
out_dims[axis] += ins[i]->dims()[j];
|
||||
continue;
|
||||
}
|
||||
PADDLE_ENFORCE_EQ(out_dims[j], ins[i]->dims()[j],
|
||||
"Input tensors should have the same "
|
||||
"elements except the specify axis.")
|
||||
}
|
||||
}
|
||||
out->Resize(out_dims);
|
||||
}
|
||||
};
|
||||
|
||||
class ConcatOpMaker : public framework::OpProtoAndCheckerMaker {
|
||||
public:
|
||||
ConcatOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
|
||||
: OpProtoAndCheckerMaker(proto, op_checker) {
|
||||
AddInput("X", "the input tensors of concat operator.").AsDuplicable();
|
||||
AddOutput("Out", "the output tensor of concat operator.");
|
||||
AddComment(R"DOC(
|
||||
Join the input tensors along with the axis.
|
||||
Examples:
|
||||
Input[0] = [[1,2],[3,4]]
|
||||
Input[1] = [[5,6]]
|
||||
axis = 0
|
||||
Output = [[1,2],
|
||||
[3,4],
|
||||
[5,6]]
|
||||
)DOC");
|
||||
AddAttr<int>("axis", "The axis which the inputs will be joined with.")
|
||||
.SetDefault(0);
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace operators
|
||||
} // namespace paddle
|
||||
|
||||
namespace ops = paddle::operators;
|
||||
REGISTER_OP_WITHOUT_GRADIENT(concat, ops::ConcatOp, ops::ConcatOpMaker)
|
||||
REGISTER_OP_CPU_KERNEL(concat,
|
||||
ops::ConcatKernel<paddle::platform::CPUPlace, float>)
|
||||
@ -0,0 +1,19 @@
|
||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License. */
|
||||
|
||||
#define EIGEN_USE_GPU
|
||||
#include "paddle/operators/concat_op.h"
|
||||
|
||||
namespace ops = paddle::operators;
|
||||
// TODO(Yancey1989) Add GPU kernel
|
||||
@ -0,0 +1,64 @@
|
||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||
|
||||
Licensed under the Apache License, Version 2.0 (the "License");
|
||||
you may not use this file except in compliance with the License.
|
||||
You may obtain a copy of the License at
|
||||
|
||||
http://www.apache.org/licenses/LICENSE-2.0
|
||||
|
||||
Unless required by applicable law or agreed to in writing, software
|
||||
distributed under the License is distributed on an "AS IS" BASIS,
|
||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
See the License for the specific language governing permissions and
|
||||
limitations under the License. */
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <vector>
|
||||
#include "paddle/framework/op_registry.h"
|
||||
|
||||
namespace paddle {
|
||||
namespace operators {
|
||||
|
||||
template <typename Place, typename T>
|
||||
class ConcatKernel : public framework::OpKernel {
|
||||
public:
|
||||
void Compute(const framework::ExecutionContext& ctx) const override {
|
||||
auto ins = ctx.MultiInput<framework::Tensor>("X");
|
||||
auto* out = ctx.Output<framework::Tensor>("Out");
|
||||
int64_t axis = static_cast<int64_t>(ctx.Attr<int>("axis"));
|
||||
size_t n = ins.size();
|
||||
size_t output_axis_dim = 0;
|
||||
size_t before = 1, after = 1;
|
||||
for (size_t i = 0; i < n; i++) {
|
||||
output_axis_dim += ins[i]->dims()[axis];
|
||||
}
|
||||
auto& input_zero = ins[0];
|
||||
for (int64_t i = 0; i < input_zero->dims().size(); i++) {
|
||||
if (i == axis) {
|
||||
continue;
|
||||
}
|
||||
if (i < axis) {
|
||||
before *= input_zero->dims()[i];
|
||||
} else {
|
||||
after *= input_zero->dims()[i];
|
||||
}
|
||||
}
|
||||
size_t output_offset = 0;
|
||||
for (size_t i = 0; i < n; i++) {
|
||||
auto& in = ins[i];
|
||||
auto axis_dim = in->dims()[axis];
|
||||
for (size_t j = 0; j < before; j++) {
|
||||
size_t len = axis_dim * after * sizeof(T);
|
||||
const T* src = in->data<T>() + axis_dim * after * j;
|
||||
T* out_data = out->mutable_data<T>(platform::CPUPlace());
|
||||
T* dest = out_data + output_offset + output_axis_dim * after * j;
|
||||
memcpy(dest, src, len);
|
||||
}
|
||||
output_offset += axis_dim * after;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace operators
|
||||
} // namespace paddle
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue