diff --git a/include/alpaka/core/Common.hpp b/include/alpaka/core/Common.hpp index a031d886ac64..5c903701a177 100644 --- a/include/alpaka/core/Common.hpp +++ b/include/alpaka/core/Common.hpp @@ -98,7 +98,7 @@ //! This macro defines a variable lying in global accelerator device memory. //! //! Example: -//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal variable; +//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal 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. @@ -114,7 +114,7 @@ //! \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 foo; +//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal foo; //! //! struct DeviceMemoryKernel //! { @@ -122,7 +122,7 @@ //! template //! ALPAKA_FN_ACC void operator()(TAcc const& acc) const //! { -//! auto a = foo.get(); +//! auto a = foo.get(); //! } //! } //! @@ -130,24 +130,24 @@ //! auto extent = alpaka::Vec, size_t>{1}; //! int initialValue = 42; //! alpaka::ViewPlainPtr, size_t> bufHost(&initialValue, devHost, extent); -//! alpaka::memcpy(queue, foo, bufHost, extent); +//! alpaka::memcpy(queue, foo, 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 \ + template \ inline __device__ #else # define ALPAKA_STATIC_ACC_MEM_GLOBAL \ - template \ + template \ inline #endif //! This macro defines a variable lying in constant accelerator device memory. //! //! Example: -//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal variable; +//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal 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. @@ -163,7 +163,7 @@ //! \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 foo; +//! ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal foo; //! //! struct DeviceMemoryKernel //! { @@ -171,7 +171,7 @@ //! template //! ALPAKA_FN_ACC void operator()(TAcc const& acc) const //! { -//! auto a = foo.get(); +//! auto a = foo.get(); //! } //! } //! @@ -179,17 +179,17 @@ //! auto extent = alpaka::Vec, size_t>{1}; //! int initialValue = 42; //! alpaka::ViewPlainPtr, size_t> bufHost(&initialValue, devHost, extent); -//! alpaka::memcpy(queue, foo, bufHost, extent); +//! alpaka::memcpy(queue, foo, 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 \ + template \ inline __constant__ #else # define ALPAKA_STATIC_ACC_MEM_CONSTANT \ - template \ + template \ inline #endif diff --git a/include/alpaka/mem/global/Traits.hpp b/include/alpaka/mem/global/Traits.hpp index 8f5c0eb5ae9a..7b3c3d1ccd3e 100644 --- a/include/alpaka/mem/global/Traits.hpp +++ b/include/alpaka/mem/global/Traits.hpp @@ -40,6 +40,6 @@ namespace alpaka }; } // namespace detail - template - using DevGlobal = typename detail::DevGlobalTrait::Type; + template + using DevGlobal = typename detail::DevGlobalTrait::type, T>::Type; } // namespace alpaka diff --git a/test/unit/mem/view/src/ViewStaticAccMem.cpp b/test/unit/mem/view/src/ViewStaticAccMem.cpp index 7ceec83614cd..a42a37c140d7 100644 --- a/test/unit/mem/view/src/ViewStaticAccMem.cpp +++ b/test/unit/mem/view/src/ViewStaticAccMem.cpp @@ -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 g_globalMemory2DUninitialized; +ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal g_globalMemory2DUninitialized; -ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal g_constantMemory2DUninitialized; +ALPAKA_STATIC_ACC_MEM_CONSTANT alpaka::DevGlobal g_constantMemory2DUninitialized; //! Uses static device memory on the accelerator defined globally for the whole compilation unit. struct StaticDeviceMemoryTestKernel @@ -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::type>.get())[0][0] - + offset)); + ALPAKA_CHECK(*success, val == *((&g_globalMemory2DUninitialized.get())[0][0] + offset)); } }; @@ -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::type>.get())[0][0] - + offset)); + ALPAKA_CHECK(*success, val == *((&g_constantMemory2DUninitialized.get())[0][0] + offset)); } }; @@ -106,22 +96,14 @@ TEMPLATE_LIST_TEST_CASE("staticDeviceMemoryGlobal", "[viewStaticAccMem]", TestAc std::vector const data{0u, 1u, 2u, 3u, 4u, 5u}; auto bufHost = alpaka::createView(devHost, data.data(), extent); - alpaka::memcpy( - queueAcc, - g_globalMemory2DUninitialized::type>, - bufHost, - extent); + alpaka::memcpy(queueAcc, g_globalMemory2DUninitialized, bufHost, extent); alpaka::wait(queueAcc); REQUIRE(fixture(kernel)); std::vector data2(6, 0u); auto bufHost2 = alpaka::createView(devHost, data2.data(), extent); - alpaka::memcpy( - queueAcc, - bufHost2, - g_globalMemory2DUninitialized::type>, - extent); + alpaka::memcpy(queueAcc, bufHost2, g_globalMemory2DUninitialized, extent); alpaka::wait(queueAcc); REQUIRE(data == data2); } @@ -152,20 +134,14 @@ TEMPLATE_LIST_TEST_CASE("staticDeviceMemoryConstant", "[viewStaticAccMem]", Test std::vector const data{0u, 1u, 2u, 3u, 4u, 5u}; auto bufHost = alpaka::createView(devHost, data.data(), extent); - alpaka::memcpy( - queueAcc, - g_constantMemory2DUninitialized::type>, - bufHost); + alpaka::memcpy(queueAcc, g_constantMemory2DUninitialized, bufHost); alpaka::wait(queueAcc); REQUIRE(fixture(kernel)); std::vector data2(6, 0u); auto bufHost2 = alpaka::createView(devHost, data2.data(), extent); - alpaka::memcpy( - queueAcc, - bufHost2, - g_constantMemory2DUninitialized::type>); + alpaka::memcpy(queueAcc, bufHost2, g_constantMemory2DUninitialized); alpaka::wait(queueAcc); REQUIRE(data == data2); }