Fixes discard_memory compilation failure for pre-Volta (#637)

* moves discard_memory to its own header

* fixes initialization order in access property

* fixes pointer alignment computation

* fixes includes and license

Author: elstehle

libcudacxx/docs/extended_api/memory_access_properties/discard_memory.md

@@ -23,7 +23,7 @@ Does **not** generate any HW store operations.
 This kernel needs a scratch pad that does not fit in shared memory, so it uses an allocation in global memory instead:
 
 ```cuda
-#include <cuda/annotated_ptr>
+#include <cuda/discard_memory>
 __device__ int compute(int* scratch, size_t N);
 
 __global__ void kernel(int const* in, int* out, int* scratch, size_t N) {

libcudacxx/include/cuda/annotated_ptr

@@ -56,6 +56,7 @@
 
 #include <cuda/std/cstdint>
 #include <cuda/barrier>
+#include <cuda/discard_memory>
 
 #include "std/detail/__access_property"
 
@@ -171,25 +172,6 @@ void apply_access_property(const volatile void* __ptr, const _Shape __shape, acc
   ))
 }
 
-inline
-_LIBCUDACXX_HOST_DEVICE
-void discard_memory(volatile void* __ptr, std::size_t __nbytes) noexcept {
-  NV_IF_TARGET(NV_PROVIDES_SM_80,(
-    if (!__isGlobal((void*)__ptr)) return;
-
-    char* __p = reinterpret_cast<char*>(const_cast<void*>(__ptr));
-    static constexpr std::size_t _LINE_SIZE = 128;
-    std::size_t __start = (reinterpret_cast<std::uintptr_t>(__p) % _LINE_SIZE) ? 1 : 0;
-    std::size_t __end = (reinterpret_cast<std::uintptr_t>(__p + __nbytes) % _LINE_SIZE) ? __nbytes - _LINE_SIZE : __nbytes;
-    __end /= _LINE_SIZE;
-
-    //Trim the first block and last block if they're not 128 bytes aligned
-    for (std::size_t __i = __start; __i < __end; __i += _LINE_SIZE) {
-      asm volatile ("discard.global.L2 [%0], 128;" ::"l"(__p + (__i * _LINE_SIZE)) :);
-    }
-  ))
-}
-
 template<class _Tp, class _Property>
 class annotated_ptr: public __detail_ap::__annotated_ptr_base<_Property> {
   public:

libcudacxx/include/cuda/discard_memory

@@ -0,0 +1,44 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_DISCARD_MEMORY
+#define _CUDA_DISCARD_MEMORY
+
+#include <cuda/std/cstdint>
+#include <cuda/std/detail/__config>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA
+
+inline _LIBCUDACXX_HOST_DEVICE void discard_memory(volatile void* __ptr, std::size_t __nbytes) noexcept
+{
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_80,
+    (if (!__isGlobal((void*) __ptr)) return;
+
+     char* __p                               = reinterpret_cast<char*>(const_cast<void*>(__ptr));
+     char* const __end_p                     = __p + __nbytes;
+     static constexpr std::size_t _LINE_SIZE = 128;
+
+     // Trim the first block and last block if they're not 128 bytes aligned
+     std::size_t __misalignment = reinterpret_cast<std::uintptr_t>(__p) % _LINE_SIZE;
+     char* __start_aligned      = __misalignment == 0 ? __p : __p + (_LINE_SIZE - __misalignment);
+     char* const __end_aligned  = __end_p - (reinterpret_cast<std::uintptr_t>(__end_p) % _LINE_SIZE);
+
+     while (__start_aligned < __end_aligned) {
+       printf("-> [%p, %p)\n", __start_aligned, __start_aligned + 128);
+       asm volatile("discard.global.L2 [%0], 128;" ::"l"(__start_aligned) :);
+       __start_aligned += _LINE_SIZE;
+     }),
+    ((void) (__ptr); (void) (__nbytes);))
+}
+
+_LIBCUDACXX_END_NAMESPACE_CUDA
+
+#endif

libcudacxx/include/cuda/std/detail/__access_property

@@ -239,13 +239,13 @@ namespace __detail_ap {
                     __on::__l2_cop_on_t __hit_prop,
                     std::uint32_t __hit_ratio,
                     __off::__l2_cop_off_t __miss_prop) noexcept
-                        : __fraction(__hit_ratio),
+                        : __ap_reserved(0x0),
+                          __fraction(__hit_ratio),
                           __l2_cop_off(__miss_prop),
                           __l2_cop_on(__hit_prop),
                           __l2_descriptor_mode(_DESC_INTERLEAVED),
                           __l1_inv_dont_allocate(0x0),
                           __l2_sector_promote_256B(0x0),
-                          __ap_reserved(0x0),
                           __ap_reserved2(0x0) {}
 
       _LIBCUDACXX_HOST_DEVICE