From 0e54517f8490816cdad1da112a41d166341088b1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 18:20:43 +0100 Subject: [PATCH] EXSWHTEC-303 - Implement tests for synchronization device functions #275 Change-Id: Iad0e303b21443615cda6fa68a9e3ef61a452a45c --- catch/include/hip_test_defgroups.hh | 7 + catch/unit/CMakeLists.txt | 1 + catch/unit/syncthreads/CMakeLists.txt | 45 ++++ catch/unit/syncthreads/__syncthreads.cc | 59 +++++ catch/unit/syncthreads/__syncthreads_and.cc | 241 ++++++++++++++++++ .../__syncthreads_and_negative_kernels.cc | 32 +++ .../__syncthreads_and_negative_kernels_rtc.hh | 39 +++ catch/unit/syncthreads/__syncthreads_count.cc | 241 ++++++++++++++++++ .../__syncthreads_count_negative_kernels.cc | 36 +++ ..._syncthreads_count_negative_kernels_rtc.hh | 39 +++ catch/unit/syncthreads/__syncthreads_or.cc | 241 ++++++++++++++++++ .../__syncthreads_or_negative_kernels.cc | 32 +++ .../__syncthreads_or_negative_kernels_rtc.hh | 39 +++ catch/unit/syncthreads/syncthreads_common.hh | 79 ++++++ 14 files changed, 1131 insertions(+) create mode 100644 catch/unit/syncthreads/CMakeLists.txt create mode 100644 catch/unit/syncthreads/__syncthreads.cc create mode 100644 catch/unit/syncthreads/__syncthreads_and.cc create mode 100644 catch/unit/syncthreads/__syncthreads_and_negative_kernels.cc create mode 100644 catch/unit/syncthreads/__syncthreads_and_negative_kernels_rtc.hh create mode 100644 catch/unit/syncthreads/__syncthreads_count.cc create mode 100644 catch/unit/syncthreads/__syncthreads_count_negative_kernels.cc create mode 100644 catch/unit/syncthreads/__syncthreads_count_negative_kernels_rtc.hh create mode 100644 catch/unit/syncthreads/__syncthreads_or.cc create mode 100644 catch/unit/syncthreads/__syncthreads_or_negative_kernels.cc create mode 100644 catch/unit/syncthreads/__syncthreads_or_negative_kernels_rtc.hh create mode 100644 catch/unit/syncthreads/syncthreads_common.hh diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 0a6d4dde90..7da19422c2 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -109,6 +109,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup SyncthreadsTest Synchronization Functions + * @{ + * This section describes tests for Synchronization Functions. + * @} + */ + /** * @defgroup MemoryTest memory Management APIs * @{ diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index f1c8faa011..3074b8b0b5 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -49,6 +49,7 @@ add_subdirectory(atomics) add_subdirectory(complex) add_subdirectory(p2p) add_subdirectory(gcc) +add_subdirectory(syncthreads) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) diff --git a/catch/unit/syncthreads/CMakeLists.txt b/catch/unit/syncthreads/CMakeLists.txt new file mode 100644 index 0000000000..0dc5faf821 --- /dev/null +++ b/catch/unit/syncthreads/CMakeLists.txt @@ -0,0 +1,45 @@ +# Copyright (c) 2023 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +set(TEST_SRC + __syncthreads.cc + __syncthreads_count.cc + __syncthreads_and.cc + __syncthreads_or.cc +) + +hip_add_exe_to_target(NAME SyncthreadsTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) + +add_test(NAME Unit___syncthreads_count_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + __syncthreads_count_negative_kernels.cc 2) + +add_test(NAME Unit___syncthreads_and_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + __syncthreads_and_negative_kernels.cc 2) + +add_test(NAME Unit___syncthreads_or_Negative_Parameters + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + __syncthreads_or_negative_kernels.cc 2) \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads.cc b/catch/unit/syncthreads/__syncthreads.cc new file mode 100644 index 0000000000..2f63ae39ae --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads.cc @@ -0,0 +1,59 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#include "syncthreads_common.hh" + +/** + * @addtogroup __syncthreads __syncthreads + * @{ + * @ingroup SyncthreadsTest + */ + +/** + * Test Description + * ------------------------ + * - Basic synchronization test for `__syncthreads`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_Positive_Basic") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, sizeof(int) * kGridSize); + + HipTest::launchKernel(SyncthreadsKernel, kGridSize, kBlockSize, + sizeof(int) * kBlockSize, nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize * (kBlockSize + 1) / 2); + } +} \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_and.cc b/catch/unit/syncthreads/__syncthreads_and.cc new file mode 100644 index 0000000000..91dc17fbb2 --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_and.cc @@ -0,0 +1,241 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#include "__syncthreads_and_negative_kernels_rtc.hh" +#include "syncthreads_common.hh" + +/** + * @addtogroup __syncthreads_and __syncthreads_and + * @{ + * @ingroup SyncthreadsTest + */ + +/** + * Test Description + * ------------------------ + * - Basic synchronization test for `__syncthreads_and`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Basic") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, sizeof(int) * kGridSize); + + HipTest::launchKernel(SyncthreadsKernel, kGridSize, kBlockSize, + sizeof(int) * kBlockSize, nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize * (kBlockSize + 1) / 2); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_and` with 0 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Predicate_Zero") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsZeroKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 0); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_and` with 1 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Predicate_One") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOneKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_and` with 0 as the predicate for even threads, and 1 as the predicate + * for odd threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Predicate_OddEven") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOddEvenKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 0); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_and` with a negative predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Predicate_Negative") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsNegativeKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_and` with the thread ID as the predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Positive_Predicate_Id") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsIdKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 0); + } +} + +/** + * Test Description + * ------------------------ + * - Real-time compiles kernels that pass invalid arguments to `__syncthreads_and`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_and.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_and_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + HIPRTC_CHECK(hiprtcCreateProgram(&program, kSyncthreadsAndSource, "__syncthreads_and_negative.cc", + 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{2}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_and_negative_kernels.cc b/catch/unit/syncthreads/__syncthreads_and_negative_kernels.cc new file mode 100644 index 0000000000..5d889a99d2 --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_and_negative_kernels.cc @@ -0,0 +1,32 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +__global__ void __syncthreads_and_v1(int* predicate) { int result = __syncthreads_and(predicate); } + +__global__ void __syncthreads_and_v2(Dummy predicate) { int result = __syncthreads_and(predicate); } \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_and_negative_kernels_rtc.hh b/catch/unit/syncthreads/__syncthreads_and_negative_kernels_rtc.hh new file mode 100644 index 0000000000..fc7be27871 --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_and_negative_kernels_rtc.hh @@ -0,0 +1,39 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kSyncthreadsAndSource{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void __syncthreads_and_v1(int* predicate) { + int result = __syncthreads_and(predicate); + } + + __global__ void __syncthreads_and_v2(Dummy predicate) { + int result = __syncthreads_and(predicate); + } + )"}; \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_count.cc b/catch/unit/syncthreads/__syncthreads_count.cc new file mode 100644 index 0000000000..dd084f436e --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_count.cc @@ -0,0 +1,241 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#include "__syncthreads_count_negative_kernels_rtc.hh" +#include "syncthreads_common.hh" + +/** + * @addtogroup __syncthreads_count __syncthreads_count + * @{ + * @ingroup SyncthreadsTest + */ + +/** + * Test Description + * ------------------------ + * - Basic synchronization test for `__syncthreads_count`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Basic") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, sizeof(int) * kGridSize); + + HipTest::launchKernel(SyncthreadsKernel, kGridSize, kBlockSize, + sizeof(int) * kBlockSize, nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize * (kBlockSize + 1) / 2); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_count` with 0 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Predicate_Zero") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsZeroKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 0); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_count` with 1 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Predicate_One") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOneKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_count` with 0 as the predicate for even threads, and 1 as the predicate + * for odd threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Predicate_OddEven") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOddEvenKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize / 2); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_count` with a negative predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Predicate_Negative") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsNegativeKernel, kGridSize, kBlockSize, + 0, nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_count` with the thread ID as the predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Positive_Predicate_Id") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsIdKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize - 1); + } +} + +/** + * Test Description + * ------------------------ + * - Real-time compiles kernels that pass invalid arguments to `__syncthreads_count`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_count.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_count_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + HIPRTC_CHECK(hiprtcCreateProgram(&program, kSyncthreadsCountSource, + "__syncthreads_count_negative.cc", 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{2}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_count_negative_kernels.cc b/catch/unit/syncthreads/__syncthreads_count_negative_kernels.cc new file mode 100644 index 0000000000..83d8cf08ab --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_count_negative_kernels.cc @@ -0,0 +1,36 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +__global__ void __syncthreads_count_v1(int* predicate) { + int result = __syncthreads_count(predicate); +} + +__global__ void __syncthreads_count_v2(Dummy predicate) { + int result = __syncthreads_count(predicate); +} \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_count_negative_kernels_rtc.hh b/catch/unit/syncthreads/__syncthreads_count_negative_kernels_rtc.hh new file mode 100644 index 0000000000..9f40e51175 --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_count_negative_kernels_rtc.hh @@ -0,0 +1,39 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kSyncthreadsCountSource{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void __syncthreads_count_v1(int* predicate) { + int result = __syncthreads_count(predicate); + } + + __global__ void __syncthreads_count_v2(Dummy predicate) { + int result = __syncthreads_count(predicate); + } + )"}; \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_or.cc b/catch/unit/syncthreads/__syncthreads_or.cc new file mode 100644 index 0000000000..d392c50eab --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_or.cc @@ -0,0 +1,241 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#include "__syncthreads_or_negative_kernels_rtc.hh" +#include "syncthreads_common.hh" + +/** + * @addtogroup __syncthreads_or __syncthreads_or + * @{ + * @ingroup SyncthreadsTest + */ + +/** + * Test Description + * ------------------------ + * - Basic synchronization test for `__syncthreads_or`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Basic") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, sizeof(int) * kGridSize); + + HipTest::launchKernel(SyncthreadsKernel, kGridSize, kBlockSize, + sizeof(int) * kBlockSize, nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == kBlockSize * (kBlockSize + 1) / 2); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_or` with 0 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Predicate_Zero") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsZeroKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 0); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_or` with 1 as the predicate for all threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Predicate_One") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOneKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_or` with 0 as the predicate for even threads, and 1 as the predicate for + * odd threads. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Predicate_OddEven") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsOddEvenKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_or` with a negative predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Predicate_Negative") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsNegativeKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Test `__syncthreads_or` with the thread ID as the predicate. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Positive_Predicate_Id") { + const auto kGridSize = 2; + const auto kBlockSize = GENERATE(13, 32, 64, 513); + + LinearAllocGuard out_alloc(LinearAllocs::hipMallocManaged, + sizeof(int) * kGridSize * kBlockSize); + + HipTest::launchKernel(SyncthreadsIdKernel, kGridSize, kBlockSize, 0, + nullptr, out_alloc.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + for (int i = 0; i < kGridSize * kBlockSize; ++i) { + REQUIRE(out_alloc.host_ptr()[i] == 1); + } +} + +/** + * Test Description + * ------------------------ + * - Real-time compiles kernels that pass invalid arguments to `__syncthreads_or`. + * + * Test source + * ------------------------ + * - unit/syncthreads/__syncthreads_or.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___syncthreads_or_Negative_Parameters_RTC") { + hiprtcProgram program{}; + + HIPRTC_CHECK(hiprtcCreateProgram(&program, kSyncthreadsOrSource, "__syncthreads_or_negative.cc", + 0, nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log and count compiler error messages + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + + int expected_error_count{2}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); + REQUIRE(error_count == expected_error_count); +} \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_or_negative_kernels.cc b/catch/unit/syncthreads/__syncthreads_or_negative_kernels.cc new file mode 100644 index 0000000000..b6f46a811c --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_or_negative_kernels.cc @@ -0,0 +1,32 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +__global__ void __syncthreads_or_v1(int* predicate) { int result = __syncthreads_or(predicate); } + +__global__ void __syncthreads_or_v2(Dummy predicate) { int result = __syncthreads_or(predicate); } \ No newline at end of file diff --git a/catch/unit/syncthreads/__syncthreads_or_negative_kernels_rtc.hh b/catch/unit/syncthreads/__syncthreads_or_negative_kernels_rtc.hh new file mode 100644 index 0000000000..dd7e1f93b0 --- /dev/null +++ b/catch/unit/syncthreads/__syncthreads_or_negative_kernels_rtc.hh @@ -0,0 +1,39 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +static constexpr auto kSyncthreadsOrSource{ + R"( + struct Dummy { + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + + __global__ void __syncthreads_or_v1(int* predicate) { + int result = __syncthreads_or(predicate); + } + + __global__ void __syncthreads_or_v2(Dummy predicate) { + int result = __syncthreads_or(predicate); + } + )"}; \ No newline at end of file diff --git a/catch/unit/syncthreads/syncthreads_common.hh b/catch/unit/syncthreads/syncthreads_common.hh new file mode 100644 index 0000000000..c6f9dec8d4 --- /dev/null +++ b/catch/unit/syncthreads/syncthreads_common.hh @@ -0,0 +1,79 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +enum class SyncthreadsKind { kDefault, kCount, kAnd, kOr }; + +template __device__ int Syncthreads(int predicate) { + if constexpr (kind == SyncthreadsKind::kDefault) { + __syncthreads(); + return 0; + } else if constexpr (kind == SyncthreadsKind::kCount) { + return __syncthreads_count(predicate); + } else if constexpr (kind == SyncthreadsKind::kAnd) { + return __syncthreads_and(predicate); + } else if constexpr (kind == SyncthreadsKind::kOr) { + return __syncthreads_or(predicate); + } +} + +template __global__ void SyncthreadsKernel(int* out) { + extern __shared__ int shared_mem[]; + + shared_mem[threadIdx.x] = threadIdx.x + 1; + + Syncthreads(0); + + if (threadIdx.x == 0) { + int sum = 0; + for (int i = 0; i < blockDim.x; ++i) { + sum += shared_mem[i]; + } + out[blockIdx.x] = sum; + } +} + +template __global__ void SyncthreadsZeroKernel(int* out) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + out[tid] = Syncthreads(0); +} + +template __global__ void SyncthreadsOneKernel(int* out) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + out[tid] = Syncthreads(1); +} + +template __global__ void SyncthreadsOddEvenKernel(int* out) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + out[tid] = Syncthreads(threadIdx.x % 2); +} + +template __global__ void SyncthreadsNegativeKernel(int* out) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + out[tid] = Syncthreads(-1); +} + +template __global__ void SyncthreadsIdKernel(int* out) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + out[tid] = Syncthreads(threadIdx.x); +} \ No newline at end of file