From 76ee8579eb1ac269da704db06dfe96039b7f45e4 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 8 Jul 2019 18:53:01 -0400 Subject: [PATCH] Check result --- tests/cuda/cuda_point.cu | 92 ++++++++++++++++++++++++++++++++-------- 1 file changed, 75 insertions(+), 17 deletions(-) diff --git a/tests/cuda/cuda_point.cu b/tests/cuda/cuda_point.cu index 15b6ba7430..e29a8f87c5 100644 --- a/tests/cuda/cuda_point.cu +++ b/tests/cuda/cuda_point.cu @@ -1,6 +1,6 @@ // --------------------------------------------------------------------- // -// Copyright (C) 2016 by the deal.II authors +// Copyright (C) 2019 by the deal.II authors // // This file is part of the deal.II library. // @@ -21,46 +21,104 @@ template __global__ void -miscellaneous_kernel() +miscellaneous_kernel(Number check[16]) { Point p_1; + check[0] = p_1.norm_square(); Point p_2(Tensor<1, dim, Number>{}); + check[1] = p_2.norm_square(); if (dim == 1) - Point p(1.); + { + Point p(1.); + check[2] = p.norm_square(); + } if (dim == 2) - Point p(1., 2.); + { + Point p(.6, .8); + check[2] = p.norm_square(); + } if (dim == 3) - Point p(1., 2., 3.); + { + Point p(.48, .64, .6); + check[2] = p.norm_square(); + } auto p_3 = Point::unit_vector(0); + check[3] = p_3.norm_square(); auto entry_1 = p_1(0); + check[4] = entry_1; p_1(0) = Number{1.}; + check[5] = p_1.norm_square(); + auto p_4 = p_1 + Tensor<1, dim, Number>{}; + check[6] = p_4.norm_square(); + auto p_5 = p_1 - Tensor<1, dim, Number>{}; + check[7] = p_5.norm_square(); + auto t_1 = p_1 - p_2; + check[8] = t_1.norm_square(); + auto p_6 = -p_3; + check[9] = p_6.norm_square(); + auto p_7 = p_4 / 2.; + check[10] = p_7.norm_square(); + auto p_8 = p_7 * 5.; + check[11] = p_8.norm_square(); - auto p_4 = p_1 + Tensor<1, dim, Number>{}; - auto p_5 = p_1 - Tensor<1, dim, Number>{}; - auto t_1 = p_1 - p_2; - auto p_6 = -p_3; - auto p_7 = p_4 / 2.; - auto p_8 = p_2 * 5.; - - auto s_1 = p_1 * t_1; - auto s_2 = p_2.square(); - auto s_3 = p_3.distance(p_5); - auto s_4 = p_4.distance_square(p_1); + auto s_1 = p_1 * t_1; + check[12] = s_1; + auto s_2 = p_2.square(); + check[13] = s_2; + auto s_3 = p_3.distance(p_5); + check[14] = s_3; + auto s_4 = p_4.distance_square(p_1); + check[15] = s_4; } template void test_gpu() { + Number * check; + const unsigned int n_tests = 16; + + auto cuda_error = cudaMalloc(&check, n_tests * sizeof(Number)); + AssertCuda(cuda_error); + // Miscellaneous - miscellaneous_kernel<<<1, 1>>>(); + miscellaneous_kernel<<<1, 1>>>(check); // Check that the kernel was launched correctly AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel AssertCuda(cudaDeviceSynchronize()); + std::vector check_host(n_tests); + + cuda_error = cudaMemcpy(check_host.data(), + check, + n_tests * sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + + const double tolerance = 1.e-8; + AssertThrow(std::abs(check_host[0] - 0.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[1] - 0.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[2] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[3] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[4] - 0.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[5] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[6] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[7] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[8] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[9] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[10] - .25) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[11] - 6.25) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[12] - 1.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[13] - 0.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[14] - 0.) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_host[15] - 0.) < tolerance, ExcInternalError()); + + cuda_error = cudaFree(check); + AssertCuda(cuda_error); + deallog << "OK" << std::endl; } -- 2.39.5