<commit_before>/* Copyright 2019 <NAME>
 *
 * This file is part of alpaka.
 *
 * This Source Code Form is subject to the terms of the Mozilla Public
 * License, v. 2.0. If a copy of the MPL was not distributed with this
 * file, You can obtain one at http://mozilla.org/MPL/2.0/.
 */

#define TEST_UNIT_KERNEL_KERNEL_STD_FUNCTION

#include <alpaka/kernel/Traits.hpp>

#include <alpaka/test/acc/TestAccs.hpp>
#include <alpaka/test/KernelExecutionFixture.hpp>
#include <alpaka/core/BoostPredef.hpp>

#include <catch2/catch.hpp>

#include <functional>
#if BOOST_LANG_CUDA
#include <nvfunctional>
#endif

//-----------------------------------------------------------------------------
template<
    typename Acc>
void ALPAKA_FN_ACC kernelFn(
    Acc const & acc,
    bool * success,
    std::int32_t val)
{
    alpaka::ignore_unused(acc);

    ALPAKA_CHECK(*success, 42 == val);
}

// std::function and std::bind is only allowed on CPU
#if !BOOST_LANG_CUDA && !BOOST_LANG_HIP
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "stdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = std::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
    REQUIRE(fixture(kernel, 42));
}

//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "stdBindKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = std::bind( kernelFn<Acc>, std::placeholders::_1, std::placeholders::_2, 42 );
    REQUIRE(fixture(kernel));
}
#endif

// This test is disabled due to #836 (cudaErrorIllegalInstruction crash)
#if 0
//#if BOOST_LANG_CUDA
// clang as a native CUDA compiler does not seem to support nvstd::function when ALPAKA_ACC_GPU_CUDA_ONLY_MODE is used.
// error: reference to __device__ function 'kernelFn<alpaka::acc::AccGpuCudaRt<std::__1::integral_constant<unsigned long, 1>, unsigned long> >' in __host__ function const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
#if !(defined(ALPAKA_ACC_GPU_CUDA_ONLY_MODE) && BOOST_COMP_CLANG_CUDA)
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "nvstdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
    REQUIRE(fixture(kernel, 42));
}

#endif
#endif<commit_msg>TEST_UNIT_KERNEL_KERNEL_STD_FUNCTION: Add clang unused macro WA<commit_after>/* Copyright 2019 <NAME>
 *
 * This file is part of alpaka.
 *
 * This Source Code Form is subject to the terms of the Mozilla Public
 * License, v. 2.0. If a copy of the MPL was not distributed with this
 * file, You can obtain one at http://mozilla.org/MPL/2.0/.
 */

#define TEST_UNIT_KERNEL_KERNEL_STD_FUNCTION
// clang thinks a macro is unused when only used as the second operand of an &&
// where the first operand evaluates to false, so we use it here:
TEST_UNIT_KERNEL_KERNEL_STD_FUNCTION

#include <alpaka/kernel/Traits.hpp>

#include <alpaka/test/acc/TestAccs.hpp>
#include <alpaka/test/KernelExecutionFixture.hpp>
#include <alpaka/core/BoostPredef.hpp>

#include <catch2/catch.hpp>

#include <functional>
#if BOOST_LANG_CUDA
#include <nvfunctional>
#endif

//-----------------------------------------------------------------------------
template<
    typename Acc>
void ALPAKA_FN_ACC kernelFn(
    Acc const & acc,
    bool * success,
    std::int32_t val)
{
    alpaka::ignore_unused(acc);

    ALPAKA_CHECK(*success, 42 == val);
}

// std::function and std::bind is only allowed on CPU
#if !BOOST_LANG_CUDA && !BOOST_LANG_HIP
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "stdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = std::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
    REQUIRE(fixture(kernel, 42));
}

//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "stdBindKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = std::bind( kernelFn<Acc>, std::placeholders::_1, std::placeholders::_2, 42 );
    REQUIRE(fixture(kernel));
}
#endif

// This test is disabled due to #836 (cudaErrorIllegalInstruction crash)
#if 0
//#if BOOST_LANG_CUDA
// clang as a native CUDA compiler does not seem to support nvstd::function when ALPAKA_ACC_GPU_CUDA_ONLY_MODE is used.
// error: reference to __device__ function 'kernelFn<alpaka::acc::AccGpuCudaRt<std::__1::integral_constant<unsigned long, 1>, unsigned long> >' in __host__ function const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
#if !(defined(ALPAKA_ACC_GPU_CUDA_ONLY_MODE) && BOOST_COMP_CLANG_CUDA)
//-----------------------------------------------------------------------------
TEMPLATE_LIST_TEST_CASE( "nvstdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
{
    using Acc = TestType;
    using Dim = alpaka::dim::Dim<Acc>;
    using Idx = alpaka::idx::Idx<Acc>;

    alpaka::test::KernelExecutionFixture<Acc> fixture(
        alpaka::vec::Vec<Dim, Idx>::ones());

    const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
    REQUIRE(fixture(kernel, 42));
}

#endif
#endif

#undef TEST_UNIT_KERNEL_KERNEL_STD_FUNCTION
{'fraction_non_alphanumeric': 0.22430769230769232, 'fraction_numerical': 0.010153846153846154, 'mean_word_length': 3.063125, 'pattern_counts': {'":': 0, '<': 65, '<?xml version=': 0, '>': 65, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2547125', 'n_tokens_mistral': 2170, 'n_tokens_neox': 2075, 'n_words': 488}