From d8fc8b2f713bbfa88a37db58376ae751e0f16226 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Mon, 8 Jul 2019 15:22:45 -0400 Subject: [PATCH] Check result --- include/deal.II/base/tensor.h | 13 ++--- tests/cuda/cuda_tensor_02.cu | 92 +++++++++++++++++++++++++++++++---- 2 files changed, 87 insertions(+), 18 deletions(-) diff --git a/include/deal.II/base/tensor.h b/include/deal.II/base/tensor.h index 5754cbc7f2..65088e7a9c 100644 --- a/include/deal.II/base/tensor.h +++ b/include/deal.II/base/tensor.h @@ -323,11 +323,9 @@ public: * Return the Frobenius-norm of a tensor, i.e. the square root of the sum of * the absolute squares of all entries. For the present case of rank-1 * tensors, this equals the usual l2 norm of the vector. - * - * @note This function can also be used in CUDA device code. */ - DEAL_II_CONSTEXPR DEAL_II_CUDA_HOST_DEV real_type - norm() const; + DEAL_II_CONSTEXPR real_type + norm() const; /** * Return the square of the Frobenius-norm of a tensor, i.e. the sum of the @@ -1036,14 +1034,11 @@ Tensor<0, dim, Number>::operator-() const template -DEAL_II_CONSTEXPR DEAL_II_CUDA_HOST_DEV inline - typename Tensor<0, dim, Number>::real_type - Tensor<0, dim, Number>::norm() const +DEAL_II_CONSTEXPR inline typename Tensor<0, dim, Number>::real_type +Tensor<0, dim, Number>::norm() const { -#ifndef __CUDA_ARCH__ Assert(dim != 0, ExcMessage("Cannot access an object of type Tensor<0,0,Number>")); -#endif return numbers::NumberTraits::abs(value); } diff --git a/tests/cuda/cuda_tensor_02.cu b/tests/cuda/cuda_tensor_02.cu index 83864d20f1..be392d5583 100644 --- a/tests/cuda/cuda_tensor_02.cu +++ b/tests/cuda/cuda_tensor_02.cu @@ -21,20 +21,29 @@ template __global__ void -miscellaneous_kernel() +miscellaneous_kernel(Number *check_1, + Number *check_2, + Number *check_3, + Number *check_4, + Number *check_5) { // constructors typename Tensor::array_type array{}; Tensor dummy_1(array); - Tensor dummy_2; - Tensor dummy_3 = dummy_2; + *check_1 = dummy_1.norm_square(); + Tensor dummy_2; + *check_2 = dummy_2.norm_square(); + Tensor dummy_3 = dummy_2; + *check_3 = dummy_3.norm_square(); // access - Tensor initializer_1; - const auto dummy_5 = initializer_1[0]; + Tensor initializer_1; + const Tensor dummy_5 = initializer_1[0]; + *check_4 = dummy_5.norm_square(); // assignment - dummy_2 = dummy_3; + dummy_2 = dummy_3; + *check_5 = dummy_2.norm_square(); } template @@ -224,9 +233,6 @@ test_gpu() AssertThrow((t2_host - reference_host).norm() < tolerance, ExcInternalError()); - // Miscellaneous - miscellaneous_kernel<<<1, 1>>>(); - // Free memory cuda_error = cudaFree(t_dev); AssertCuda(cuda_error); @@ -234,6 +240,74 @@ test_gpu() AssertCuda(cuda_error); cuda_error = cudaFree(t2_dev); AssertCuda(cuda_error); + + // Miscellaneous + { + Number *check_1; + Number *check_2; + Number *check_3; + Number *check_4; + Number *check_5; + + cuda_error = cudaMalloc(&check_1, sizeof(Number)); + AssertCuda(cuda_error); + cuda_error = cudaMalloc(&check_2, sizeof(Number)); + AssertCuda(cuda_error); + cuda_error = cudaMalloc(&check_3, sizeof(Number)); + AssertCuda(cuda_error); + cuda_error = cudaMalloc(&check_4, sizeof(Number)); + AssertCuda(cuda_error); + cuda_error = cudaMalloc(&check_5, sizeof(Number)); + AssertCuda(cuda_error); + + miscellaneous_kernel + <<<1, 1>>>(check_1, check_2, check_3, check_4, check_5); + + Number check_1_host, check_2_host, check_3_host, check_4_host, check_5_host; + + cuda_error = cudaMemcpy(&check_1_host, + check_1, + sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + cuda_error = cudaMemcpy(&check_2_host, + check_2, + sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + cuda_error = cudaMemcpy(&check_3_host, + check_3, + sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + cuda_error = cudaMemcpy(&check_4_host, + check_4, + sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + cuda_error = cudaMemcpy(&check_5_host, + check_5, + sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(cuda_error); + + AssertThrow(std::abs(check_1_host) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_2_host) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_3_host) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_4_host) < tolerance, ExcInternalError()); + AssertThrow(std::abs(check_5_host) < tolerance, ExcInternalError()); + + cuda_error = cudaFree(check_1); + AssertCuda(cuda_error); + cuda_error = cudaFree(check_2); + AssertCuda(cuda_error); + cuda_error = cudaFree(check_3); + AssertCuda(cuda_error); + cuda_error = cudaFree(check_4); + AssertCuda(cuda_error); + cuda_error = cudaFree(check_5); + AssertCuda(cuda_error); + } } int -- 2.39.5