aboutsummaryrefslogtreecommitdiffhomepage
path: root/unsupported/test/cxx11_tensor_device.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'unsupported/test/cxx11_tensor_device.cpp')
-rw-r--r--unsupported/test/cxx11_tensor_device.cpp118
1 files changed, 59 insertions, 59 deletions
diff --git a/unsupported/test/cxx11_tensor_device.cpp b/unsupported/test/cxx11_tensor_device.cpp
index 26465ee11..f2d7e4ce6 100644
--- a/unsupported/test/cxx11_tensor_device.cpp
+++ b/unsupported/test/cxx11_tensor_device.cpp
@@ -22,23 +22,23 @@ using Eigen::RowMajor;
// Context for evaluation on cpu
struct CPUContext {
- CPUContext(const Eigen::Tensor<float, 3>& in1, Eigen::Tensor<float, 3>& in2, Eigen::Tensor<float, 3>& out) : in1_(in1), in2_(in2), out_(out), kernel_1d_(2), kernel_2d_(Eigen::array<int, 2>(2,2)), kernel_3d_(Eigen::array<int, 3>(2,2,2)) {
+ CPUContext(const Eigen::Tensor<float, 3>& in1, Eigen::Tensor<float, 3>& in2, Eigen::Tensor<float, 3>& out) : in1_(in1), in2_(in2), out_(out), kernel_1d_(2), kernel_2d_(2,2), kernel_3d_(2,2,2) {
kernel_1d_(0) = 3.14f;
kernel_1d_(1) = 2.7f;
- kernel_2d_(Eigen::array<int, 2>(0,0)) = 3.14f;
- kernel_2d_(Eigen::array<int, 2>(1,0)) = 2.7f;
- kernel_2d_(Eigen::array<int, 2>(0,1)) = 0.2f;
- kernel_2d_(Eigen::array<int, 2>(1,1)) = 7.0f;
-
- kernel_3d_(Eigen::array<int, 3>(0,0,0)) = 3.14f;
- kernel_3d_(Eigen::array<int, 3>(0,1,0)) = 2.7f;
- kernel_3d_(Eigen::array<int, 3>(0,0,1)) = 0.2f;
- kernel_3d_(Eigen::array<int, 3>(0,1,1)) = 7.0f;
- kernel_3d_(Eigen::array<int, 3>(1,0,0)) = -1.0f;
- kernel_3d_(Eigen::array<int, 3>(1,1,0)) = -0.3f;
- kernel_3d_(Eigen::array<int, 3>(1,0,1)) = -0.7f;
- kernel_3d_(Eigen::array<int, 3>(1,1,1)) = -0.5f;
+ kernel_2d_(0,0) = 3.14f;
+ kernel_2d_(1,0) = 2.7f;
+ kernel_2d_(0,1) = 0.2f;
+ kernel_2d_(1,1) = 7.0f;
+
+ kernel_3d_(0,0,0) = 3.14f;
+ kernel_3d_(0,1,0) = 2.7f;
+ kernel_3d_(0,0,1) = 0.2f;
+ kernel_3d_(0,1,1) = 7.0f;
+ kernel_3d_(1,0,0) = -1.0f;
+ kernel_3d_(1,1,0) = -0.3f;
+ kernel_3d_(1,0,1) = -0.7f;
+ kernel_3d_(1,1,1) = -0.5f;
}
const Eigen::DefaultDevice& device() const { return cpu_device_; }
@@ -93,8 +93,8 @@ struct GPUContext {
const Eigen::TensorMap<Eigen::Tensor<float, 3> >& in2() const { return in2_; }
Eigen::TensorMap<Eigen::Tensor<float, 3> >& out() { return out_; }
Eigen::TensorMap<Eigen::Tensor<float, 1> > kernel1d() const { return Eigen::TensorMap<Eigen::Tensor<float, 1> >(kernel_1d_, 2); }
- Eigen::TensorMap<Eigen::Tensor<float, 2> > kernel2d() const { return Eigen::TensorMap<Eigen::Tensor<float, 2> >(kernel_2d_, Eigen::array<int, 2>(2, 2)); }
- Eigen::TensorMap<Eigen::Tensor<float, 3> > kernel3d() const { return Eigen::TensorMap<Eigen::Tensor<float, 3> >(kernel_3d_, Eigen::array<int, 3>(2, 2, 2)); }
+ Eigen::TensorMap<Eigen::Tensor<float, 2> > kernel2d() const { return Eigen::TensorMap<Eigen::Tensor<float, 2> >(kernel_2d_, 2, 2); }
+ Eigen::TensorMap<Eigen::Tensor<float, 3> > kernel3d() const { return Eigen::TensorMap<Eigen::Tensor<float, 3> >(kernel_3d_, 2, 2, 2); }
private:
const Eigen::TensorMap<Eigen::Tensor<float, 3> >& in1_;
@@ -150,8 +150,8 @@ static void test_contraction(Context* context)
template <typename Context>
static void test_1d_convolution(Context* context)
{
- Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0));
- Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,70));
+ Eigen::DSizes<int, 3> indices(0,0,0);
+ Eigen::DSizes<int, 3> sizes(40,49,70);
Eigen::array<int, 1> dims(1);
context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel1d(), dims);
@@ -160,8 +160,8 @@ static void test_1d_convolution(Context* context)
template <typename Context>
static void test_2d_convolution(Context* context)
{
- Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0));
- Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,69));
+ Eigen::DSizes<int, 3> indices(0,0,0);
+ Eigen::DSizes<int, 3> sizes(40,49,69);
Eigen::array<int, 2> dims(1,2);
context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel2d(), dims);
@@ -170,8 +170,8 @@ static void test_2d_convolution(Context* context)
template <typename Context>
static void test_3d_convolution(Context* context)
{
- Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0));
- Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(39,49,69));
+ Eigen::DSizes<int, 3> indices(0,0,0);
+ Eigen::DSizes<int, 3> sizes(39,49,69);
Eigen::array<int, 3> dims(0,1,2);
context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel3d(), dims);
@@ -179,9 +179,9 @@ static void test_3d_convolution(Context* context)
static void test_cpu() {
- Eigen::Tensor<float, 3> in1(Eigen::array<int, 3>(40,50,70));
- Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70));
- Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(40,50,70));
+ Eigen::Tensor<float, 3> in1(40,50,70);
+ Eigen::Tensor<float, 3> in2(40,50,70);
+ Eigen::Tensor<float, 3> out(40,50,70);
in1 = in1.random() + in1.constant(10.0f);
in2 = in2.random() + in2.constant(10.0f);
@@ -191,7 +191,7 @@ static void test_cpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f);
}
}
}
@@ -200,7 +200,7 @@ static void test_cpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k))) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) + in2(i,j,k)) * 3.14f + 2.718f);
}
}
}
@@ -209,7 +209,7 @@ static void test_cpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f);
}
}
}
@@ -217,11 +217,11 @@ static void test_cpu() {
test_contraction(&context);
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 40; ++j) {
- const float result = out(Eigen::array<int, 3>(i,j,0));
+ const float result = out(i,j,0);
float expected = 0;
for (int k = 0; k < 50; ++k) {
for (int l = 0; l < 70; ++l) {
- expected += in1(Eigen::array<int, 3>(i, k, l)) * in2(Eigen::array<int, 3>(j, k, l));
+ expected += in1(i, k, l) * in2(j, k, l);
}
}
VERIFY_IS_APPROX(expected, result);
@@ -232,7 +232,7 @@ static void test_cpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f));
+ VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f));
}
}
}
@@ -241,9 +241,9 @@ static void test_cpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 69; ++k) {
- const float result = out(Eigen::array<int, 3>(i,j,k));
- const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f) +
- (in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f);
+ const float result = out(i,j,k);
+ const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f) +
+ (in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f);
if (fabs(expected) < 1e-4 && fabs(result) < 1e-4) {
continue;
}
@@ -256,11 +256,11 @@ static void test_cpu() {
for (int i = 0; i < 39; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 69; ++k) {
- const float result = out(Eigen::array<int, 3>(i,j,k));
- const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f +
- in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f) +
- (in1(Eigen::array<int, 3>(i+1,j,k)) * -1.0f + in1(Eigen::array<int, 3>(i+1,j+1,k)) * -0.3f +
- in1(Eigen::array<int, 3>(i+1,j,k+1)) * -0.7f + in1(Eigen::array<int, 3>(i+1,j+1,k+1)) * -0.5f);
+ const float result = out(i,j,k);
+ const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f +
+ in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f) +
+ (in1(i+1,j,k) * -1.0f + in1(i+1,j+1,k) * -0.3f +
+ in1(i+1,j,k+1) * -0.7f + in1(i+1,j+1,k+1) * -0.5f);
if (fabs(expected) < 1e-4 && fabs(result) < 1e-4) {
continue;
}
@@ -271,9 +271,9 @@ static void test_cpu() {
}
static void test_gpu() {
- Eigen::Tensor<float, 3> in1(Eigen::array<int, 3>(40,50,70));
- Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70));
- Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(40,50,70));
+ Eigen::Tensor<float, 3> in1(40,50,70);
+ Eigen::Tensor<float, 3> in2(40,50,70);
+ Eigen::Tensor<float, 3> out(40,50,70);
in1 = in1.random() + in1.constant(10.0f);
in2 = in2.random() + in2.constant(10.0f);
@@ -291,9 +291,9 @@ static void test_gpu() {
cudaMemcpy(d_in1, in1.data(), in1_bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_in2, in2.data(), in2_bytes, cudaMemcpyHostToDevice);
- Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in1(d_in1, Eigen::array<int, 3>(40,50,70));
- Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in2(d_in2, Eigen::array<int, 3>(40,50,70));
- Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_out(d_out, Eigen::array<int, 3>(40,50,70));
+ Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in1(d_in1, 40,50,70);
+ Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in2(d_in2, 40,50,70);
+ Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_out(d_out, 40,50,70);
GPUContext context(gpu_in1, gpu_in2, gpu_out);
test_contextual_eval(&context);
@@ -301,7 +301,7 @@ static void test_gpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f);
}
}
}
@@ -311,7 +311,7 @@ static void test_gpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k))) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) + in2(i,j,k)) * 3.14f + 2.718f);
}
}
}
@@ -321,7 +321,7 @@ static void test_gpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 50; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f);
+ VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f);
}
}
}
@@ -330,11 +330,11 @@ static void test_gpu() {
assert(cudaMemcpy(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost) == cudaSuccess);
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 40; ++j) {
- const float result = out(Eigen::array<int, 3>(i,j,0));
+ const float result = out(i,j,0);
float expected = 0;
for (int k = 0; k < 50; ++k) {
for (int l = 0; l < 70; ++l) {
- expected += in1(Eigen::array<int, 3>(i, k, l)) * in2(Eigen::array<int, 3>(j, k, l));
+ expected += in1(i, k, l) * in2(j, k, l);
}
}
VERIFY_IS_APPROX(expected, result);
@@ -347,7 +347,7 @@ static void test_gpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 70; ++k) {
- VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f));
+ VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f));
}
}
}
@@ -358,9 +358,9 @@ static void test_gpu() {
for (int i = 0; i < 40; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 69; ++k) {
- const float result = out(Eigen::array<int, 3>(i,j,k));
- const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f +
- in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f);
+ const float result = out(i,j,k);
+ const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f +
+ in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f);
VERIFY_IS_APPROX(expected, result);
}
}
@@ -372,11 +372,11 @@ static void test_gpu() {
for (int i = 0; i < 39; ++i) {
for (int j = 0; j < 49; ++j) {
for (int k = 0; k < 69; ++k) {
- const float result = out(Eigen::array<int, 3>(i,j,k));
- const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f +
- in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f +
- in1(Eigen::array<int, 3>(i+1,j,k)) * -1.0f + in1(Eigen::array<int, 3>(i+1,j+1,k)) * -0.3f +
- in1(Eigen::array<int, 3>(i+1,j,k+1)) * -0.7f + in1(Eigen::array<int, 3>(i+1,j+1,k+1)) * -0.5f);
+ const float result = out(i,j,k);
+ const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f +
+ in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f +
+ in1(i+1,j,k) * -1.0f + in1(i+1,j+1,k) * -0.3f +
+ in1(i+1,j,k+1) * -0.7f + in1(i+1,j+1,k+1) * -0.5f);
VERIFY_IS_APPROX(expected, result);
}
}
generated by cgit v1.2.3 (git 2.39.1) at 2024-06-06 01:34:29 +0000