Skip to content

Commit

Permalink
move Tag inside DevGlobalTrait, the user will use Acc
Browse files Browse the repository at this point in the history
  • Loading branch information
@AuroraPerego @psychocoderHPC
AuroraPerego authored and psychocoderHPC committed Apr 16, 2024
1 parent a0c53f0 commit 88c776b
Show file tree
Hide file tree
Showing 3 changed files with 22 additions and 46 deletions.
24 changes: 12 additions & 12 deletions include/alpaka/core/Common.hpp
View file
Expand Up @@ -98,7 +98,7 @@
//! This macro defines a variable lying in global accelerator device memory.
//!
//! Example:
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<int> variable;
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TAcc, int> variable;
//!
//! Those variables behave like ordinary variables when used in file-scope,
//! but inside kernels the get() method must be used to access the variable.
Expand All @@ -114,40 +114,40 @@
//! \attention It is not allowed to initialize the variable together with the declaration.
//! To initialize the variable alpaka::memcpy must be used.
//! \code{.cpp}
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<int> foo;
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TAcc, int> foo;
//!
//! struct DeviceMemoryKernel
//! {
//! ALPAKA_NO_HOST_ACC_WARNING
//! template<typename TAcc>
//! ALPAKA_FN_ACC void operator()(TAcc const& acc) const
//! {
//! auto a = foo<Tag>.get();
//! auto a = foo<TAcc>.get();
//! }
//! }
//!
//! void initFoo() {
//! auto extent = alpaka::Vec<alpaka::DimInt<1u>, size_t>{1};
//! int initialValue = 42;
//! alpaka::ViewPlainPtr<DevHost, int, alpaka::DimInt<1u>, size_t> bufHost(&initialValue, devHost, extent);
//! alpaka::memcpy(queue, foo<Tag>, bufHost, extent);
//! alpaka::memcpy(queue, foo<Acc>, bufHost, extent);
//! }
//! \endcode
#if((BOOST_LANG_CUDA && BOOST_COMP_CLANG_CUDA) || (BOOST_LANG_CUDA && BOOST_COMP_NVCC && BOOST_ARCH_PTX) \
|| BOOST_LANG_HIP)
# define ALPAKA_STATIC_ACC_MEM_GLOBAL \
template<typename TTag> \
template<typename TAcc> \
inline __device__
#else
# define ALPAKA_STATIC_ACC_MEM_GLOBAL \
template<typename TTag> \
template<typename TAcc> \
inline
#endif

//! This macro defines a variable lying in constant accelerator device memory.
//!
//! Example:
//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<int> variable;
//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<TAcc, const int> variable;
//!
//! Those variables behave like ordinary variables when used in file-scope,
//! but inside kernels the get() method must be used to access the variable.
Expand All @@ -163,33 +163,33 @@
//! \attention It is not allowed to initialize the variable together with the declaration.
//! To initialize the variable alpaka::memcpy must be used.
//! \code{.cpp}
//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<int> foo;
//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<TAcc, const int> foo;
//!
//! struct DeviceMemoryKernel
//! {
//! ALPAKA_NO_HOST_ACC_WARNING
//! template<typename TAcc>
//! ALPAKA_FN_ACC void operator()(TAcc const& acc) const
//! {
//! auto a = foo<Tag>.get();
//! auto a = foo<TAcc>.get();
//! }
//! }
//!
//! void initFoo() {
//! auto extent = alpaka::Vec<alpaka::DimInt<1u>, size_t>{1};
//! int initialValue = 42;
//! alpaka::ViewPlainPtr<DevHost, int, alpaka::DimInt<1u>, size_t> bufHost(&initialValue, devHost, extent);
//! alpaka::memcpy(queue, foo<Tag>, bufHost, extent);
//! alpaka::memcpy(queue, foo<Acc>, bufHost, extent);
//! }
//! \endcode
#if((BOOST_LANG_CUDA && BOOST_COMP_CLANG_CUDA) || (BOOST_LANG_CUDA && BOOST_COMP_NVCC && BOOST_ARCH_PTX) \
|| BOOST_LANG_HIP)
# define ALPAKA_STATIC_ACC_MEM_CONSTANT \
template<typename TTag> \
template<typename TAcc> \
inline __constant__
#else
# define ALPAKA_STATIC_ACC_MEM_CONSTANT \
template<typename TTag> \
template<typename TAcc> \
inline
#endif

Expand Down
4 changes: 2 additions & 2 deletions include/alpaka/mem/global/Traits.hpp
View file
Expand Up @@ -40,6 +40,6 @@ namespace alpaka
};
} // namespace detail

template<typename TTag, typename T>
using DevGlobal = typename detail::DevGlobalTrait<TTag, T>::Type;
template<typename TAcc, typename T>
using DevGlobal = typename detail::DevGlobalTrait<typename alpaka::trait::AccToTag<TAcc>::type, T>::Type;
} // namespace alpaka
40 changes: 8 additions & 32 deletions test/unit/mem/view/src/ViewStaticAccMem.cpp
View file
Expand Up @@ -16,9 +16,9 @@ using Elem = std::uint32_t;
using Dim = alpaka::DimInt<2u>;
using Idx = std::uint32_t;

ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TTag, Elem[3][2]> g_globalMemory2DUninitialized;
ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TAcc, Elem[3][2]> g_globalMemory2DUninitialized;

ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<TTag, const Elem[3][2]> g_constantMemory2DUninitialized;
ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal<TAcc, const Elem[3][2]> g_constantMemory2DUninitialized;

//! Uses static device memory on the accelerator defined globally for the whole compilation unit.
struct StaticDeviceMemoryTestKernel
Expand All @@ -34,12 +34,7 @@ struct StaticDeviceMemoryTestKernel
auto const offset = gridThreadExtent[1u] * gridThreadIdx[0u] + gridThreadIdx[1u];
auto const val = offset;

ALPAKA_CHECK(
*success,
val
== *(
(&g_globalMemory2DUninitialized<typename alpaka::trait::AccToTag<TAcc>::type>.get())[0][0]
+ offset));
ALPAKA_CHECK(*success, val == *((&g_globalMemory2DUninitialized<TAcc>.get())[0][0] + offset));
}
};

Expand All @@ -55,12 +50,7 @@ struct ConstantDeviceMemoryTestKernel
auto const offset = gridThreadExtent[1u] * gridThreadIdx[0u] + gridThreadIdx[1u];
auto const val = offset;

ALPAKA_CHECK(
*success,
val
== *(
(&g_constantMemory2DUninitialized<typename alpaka::trait::AccToTag<TAcc>::type>.get())[0][0]
+ offset));
ALPAKA_CHECK(*success, val == *((&g_constantMemory2DUninitialized<TAcc>.get())[0][0] + offset));
}
};

Expand Down Expand Up @@ -106,22 +96,14 @@ TEMPLATE_LIST_TEST_CASE("staticDeviceMemoryGlobal", "[viewStaticAccMem]", TestAc
std::vector<Elem> const data{0u, 1u, 2u, 3u, 4u, 5u};
auto bufHost = alpaka::createView(devHost, data.data(), extent);

alpaka::memcpy(
queueAcc,
g_globalMemory2DUninitialized<typename alpaka::trait::AccToTag<Acc>::type>,
bufHost,
extent);
alpaka::memcpy(queueAcc, g_globalMemory2DUninitialized<Acc>, bufHost, extent);
alpaka::wait(queueAcc);

REQUIRE(fixture(kernel));

std::vector<Elem> data2(6, 0u);
auto bufHost2 = alpaka::createView(devHost, data2.data(), extent);
alpaka::memcpy(
queueAcc,
bufHost2,
g_globalMemory2DUninitialized<typename alpaka::trait::AccToTag<Acc>::type>,
extent);
alpaka::memcpy(queueAcc, bufHost2, g_globalMemory2DUninitialized<Acc>, extent);
alpaka::wait(queueAcc);
REQUIRE(data == data2);
}
Expand Down Expand Up @@ -152,20 +134,14 @@ TEMPLATE_LIST_TEST_CASE("staticDeviceMemoryConstant", "[viewStaticAccMem]", Test
std::vector<Elem> const data{0u, 1u, 2u, 3u, 4u, 5u};
auto bufHost = alpaka::createView(devHost, data.data(), extent);

alpaka::memcpy(
queueAcc,
g_constantMemory2DUninitialized<typename alpaka::trait::AccToTag<Acc>::type>,
bufHost);
alpaka::memcpy(queueAcc, g_constantMemory2DUninitialized<Acc>, bufHost);
alpaka::wait(queueAcc);

REQUIRE(fixture(kernel));

std::vector<Elem> data2(6, 0u);
auto bufHost2 = alpaka::createView(devHost, data2.data(), extent);
alpaka::memcpy(
queueAcc,
bufHost2,
g_constantMemory2DUninitialized<typename alpaka::trait::AccToTag<Acc>::type>);
alpaka::memcpy(queueAcc, bufHost2, g_constantMemory2DUninitialized<Acc>);
alpaka::wait(queueAcc);
REQUIRE(data == data2);
}
Expand Down

0 comments on commit 88c776b

Please sign in to comment.