From ecf57ee0d8d31d4319eb6bbcef8eade94e73d574 Mon Sep 17 00:00:00 2001 From: Kjetil Olsen Lye Date: Thu, 6 Feb 2025 23:15:48 +0100 Subject: [PATCH] Added tests for error macros on GPU. --- CMakeLists.txt | 1 + CMakeLists_files.cmake | 2 + tests/gpuistl/test_throw_macros_on_gpu.cu | 109 ++++++++++++++++++++++ 3 files changed, 112 insertions(+) create mode 100644 tests/gpuistl/test_throw_macros_on_gpu.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index a40adc97e..2339667d8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -731,6 +731,7 @@ if(CUDA_FOUND) gpu_resources gpu_smart_pointers is_gpu_pointer + throw_macros_on_gpu PROPERTIES LABELS ${gpu_label}) endif() diff --git a/CMakeLists_files.cmake b/CMakeLists_files.cmake index 9d9dded87..a11d14c37 100644 --- a/CMakeLists_files.cmake +++ b/CMakeLists_files.cmake @@ -480,6 +480,8 @@ if (HAVE_CUDA) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_smart_pointers.cu) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_gpu_resources.cu) ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_is_gpu_pointer.cpp) + ADD_CUDA_OR_HIP_FILE(TEST_SOURCE_FILES tests test_throw_macros_on_gpu.cu) + # for loop providing the flag --expt-relaxed-constexpr to fix some cuda issues with constexpr if(NOT CONVERT_CUDA_TO_HIP) diff --git a/tests/gpuistl/test_throw_macros_on_gpu.cu b/tests/gpuistl/test_throw_macros_on_gpu.cu new file mode 100644 index 000000000..800ed6149 --- /dev/null +++ b/tests/gpuistl/test_throw_macros_on_gpu.cu @@ -0,0 +1,109 @@ +/* + Copyright 2025 Equinor ASA + + This file is part of the Open Porous Media project (OPM). + OPM is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + OPM is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + You should have received a copy of the GNU General Public License + along with OPM. If not, see . +*/ +#include +#include +#include + +#define BOOST_TEST_MODULE TestThrowMacrosOnGPU + +#include +#include +#include +#include +#include + +namespace { +__global__ void codeThatContainsMacros(bool call) { + if (call) { + OPM_THROW(std::logic_error, "Something went wrong"); + OPM_THROW_NOLOG(std::logic_error, "Something went wrong"); + OPM_THROW_PROBLEM(std::logic_error, "Something went wrong"); + } + OPM_ERROR_IF(!call, "Something went horribly wrong"); +} + +// TODO: Check if this is better on HIP +#if 0 // I am leaving this here to show that this is not possible due to limitations in CUDA + // the assert will indeed cause an error, but the CUDA context will be broken for + // the rest of the lifetime of the process, see + // https://forums.developer.nvidia.com/t/how-to-clear-cuda-errors/296393/5 +__global__ void checkThrow() { + OPM_THROW(std::logic_error, "Something went wrong"); +} + +__global__ void checkThrowNoLog() { + OPM_THROW_NOLOG(std::logic_error, "Something went wrong"); +} + +__global__ void checkThrowProblem() { + OPM_THROW_PROBLEM(std::logic_error, "Something went wrong"); +} + +__global__ void checkErrorIf() { + OPM_ERROR_IF(true, "Something went horribly wrong"); +} +#endif +} + +BOOST_AUTO_TEST_CASE(TestKernel) +{ + OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + codeThatContainsMacros<<<1, 1>>>(false); + OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + + #if 0 // I am leaving this here to show that this is not possible due to limitations in CUDA + // the assert will indeed cause an error, but the CUDA context will be broken for + // the rest of the lifetime of the process, see + // https://forums.developer.nvidia.com/t/how-to-clear-cuda-errors/296393/5 + codeThatContainsMacros<<<1, 1>>>(true); + // Make sure this actually throws + BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error); + OPM_GPU_SAFE_CALL(cudaDeviceReset()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + + checkThrow<<<1, 1>>>(); + BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error); + OPM_GPU_SAFE_CALL(cudaDeviceReset()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + + checkThrowNoLog<<<1, 1>>>(); + BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error); + OPM_GPU_SAFE_CALL(cudaDeviceReset()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + + checkThrowProblem<<<1, 1>>>(); + BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error); + OPM_GPU_SAFE_CALL(cudaDeviceReset()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + + checkErrorIf<<<1, 1>>>(); + BOOST_CHECK_THROW(OPM_GPU_SAFE_CALL(cudaDeviceSynchronize()), std::runtime_error); + OPM_GPU_SAFE_CALL(cudaDeviceReset()); + OPM_GPU_SAFE_CALL(cudaGetLastError()); + #endif +} + +BOOST_AUTO_TEST_CASE(TestOutsideKernel) +{ + // This is to make sure that the macros work outside of kernels but inside a .cu file + // ie. inside a file compiled by nvcc/hipcc. + BOOST_CHECK_THROW(OPM_THROW(std::runtime_error, "THROW"), std::runtime_error); + BOOST_CHECK_THROW(OPM_THROW_NOLOG(std::runtime_error, "THROW_NOLOG"), std::runtime_error); + BOOST_CHECK_THROW(OPM_THROW_PROBLEM(std::runtime_error, "THROW_PROBLEM"), std::runtime_error); + BOOST_CHECK_THROW(OPM_ERROR_IF(true, "ERROR_IF"), std::logic_error); +}