diff --git a/include/alpaka/alpaka.hpp b/include/alpaka/alpaka.hpp
index f1f271833da1..d82ddfe8adcb 100644
--- a/include/alpaka/alpaka.hpp
+++ b/include/alpaka/alpaka.hpp
@@ -68,6 +68,7 @@
 #include <alpaka/core/Hip.hpp>
 #include <alpaka/core/OmpSchedule.hpp>
 #include <alpaka/core/Positioning.hpp>
+#include <alpaka/core/RemoveRestrict.hpp>
 #include <alpaka/core/Unroll.hpp>
 #include <alpaka/core/Unused.hpp>
 #include <alpaka/core/Utility.hpp>
diff --git a/include/alpaka/core/RemoveRestrict.hpp b/include/alpaka/core/RemoveRestrict.hpp
new file mode 100644
index 000000000000..c87af8729042
--- /dev/null
+++ b/include/alpaka/core/RemoveRestrict.hpp
@@ -0,0 +1,40 @@
+/* Copyright 2021 Rene Widera
+ *
+ * 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/.
+ */
+
+#pragma once
+
+#include <alpaka/core/BoostPredef.hpp>
+
+namespace alpaka
+{
+    //! Removes __restrict__ from a type
+    template<typename T>
+    struct remove_restrict
+    {
+        using type = T;
+    };
+
+#if BOOST_COMP_MSVC
+    template<typename T>
+    struct remove_restrict<T __restrict>
+    {
+        using type = T;
+    };
+#else
+    template<typename T>
+    struct remove_restrict<T __restrict__>
+    {
+        using type = T;
+    };
+#endif
+
+    //! Helper to remove __restrict__ from a type
+    template<typename T>
+    using remove_restrict_t = typename remove_restrict<T>::type;
+} // namespace alpaka
diff --git a/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp b/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp
index 416bceacf429..f5d4a1359def 100644
--- a/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp
+++ b/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp
@@ -39,6 +39,7 @@
 // Implementation details.
 #    include <alpaka/acc/AccGpuUniformCudaHipRt.hpp>
 #    include <alpaka/core/Decay.hpp>
+#    include <alpaka/core/RemoveRestrict.hpp>
 #    include <alpaka/core/Unused.hpp>
 #    include <alpaka/dev/DevUniformCudaHipRt.hpp>
 #    include <alpaka/kernel/Traits.hpp>
@@ -75,7 +76,7 @@ namespace alpaka
             __global__ void uniformCudaHipKernel(
                 Vec<TDim, TIdx> const threadElemExtent,
                 TKernelFnObj const kernelFnObj,
-                TArgs... args)
+                remove_restrict_t<TArgs>... args)
             {
 #    if BOOST_ARCH_PTX && (BOOST_ARCH_PTX < BOOST_VERSION_NUMBER(2, 0, 0))
 #        error "Device capability >= 2.0 is required!"