Commit 4d1a922e authored by carlushuang's avatar carlushuang
Browse files

control memory coherence while construct dynamic buffer

parent 5d11ef6d
......@@ -19,7 +19,8 @@ namespace ck {
template <AddressSpaceEnum BufferAddressSpace,
typename T,
typename ElementSpaceSize,
bool InvalidElementUseNumericalZeroValue>
bool InvalidElementUseNumericalZeroValue,
amd_buffer_coherence_bits coherence = amd_buffer_coherence_bits::default_coherence>
struct DynamicBuffer
{
using type = T;
......@@ -52,10 +53,9 @@ struct DynamicBuffer
__host__ __device__ constexpr T& operator()(index_t i) { return p_data_[i]; }
template <typename X,
amd_buffer_coherence_bits coherence = amd_buffer_coherence_bits::default_coherence,
typename enable_if<is_same<typename scalar_type<remove_cvref_t<X>>::type,
typename scalar_type<remove_cvref_t<T>>::type>::value,
bool>::type = false>
bool>::type = false>
__host__ __device__ constexpr auto Get(index_t i, bool is_valid_element) const
{
// X contains multiple T
......@@ -148,10 +148,9 @@ struct DynamicBuffer
}
template <typename X,
amd_buffer_coherence_bits coherence = amd_buffer_coherence_bits::default_coherence,
typename enable_if<is_same<typename scalar_type<remove_cvref_t<X>>::type,
typename scalar_type<remove_cvref_t<T>>::type>::value,
bool>::type = false>
bool>::type = false>
__host__ __device__ void Set(index_t i, bool is_valid_element, const X& x)
{
// X contains multiple T
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment