Use DynamicBuffer instead of raw pointer (#32)

* Use DynamicBuffer to hold raw pointer (to global and LDS memory)

* add workaround for compiler issue (inefficient ISA) of ds_write for int8x4, int8x8, int8x16

[ROCm/composable_kernel commit: 78b987fbd6]

composable_kernel/include/utility/static_buffer.hpp

@@ -0,0 +1,33 @@
+#ifndef CK_STATIC_BUFFER_HPP
+#define CK_STATIC_BUFFER_HPP
+
+#include "statically_indexed_array.hpp"
+
+namespace ck {
+
+template <AddressSpace BufferAddressSpace, typename T, index_t N>
+struct StaticBuffer : public StaticallyIndexedArray<T, N>
+{
+    using type = T;
+    using base = StaticallyIndexedArray<T, N>;
+
+    __host__ __device__ constexpr StaticBuffer() : base{} {}
+
+    __host__ __device__ static constexpr AddressSpace GetAddressSpace()
+    {
+        return BufferAddressSpace;
+    }
+
+    __host__ __device__ static constexpr bool IsStaticBuffer() { return true; }
+
+    __host__ __device__ static constexpr bool IsDynamicBuffer() { return false; }
+};
+
+template <AddressSpace BufferAddressSpace = AddressSpace::Generic, typename T, index_t N>
+__host__ __device__ constexpr auto make_static_buffer(Number<N>)
+{
+    return StaticBuffer<BufferAddressSpace, T, N>{};
+}
+
+} // namespace ck
+#endif