Skip to content

GitLab

"tests/pipelines/vscode:/vscode.git/clone" did not exist on "fd3d5502d4ab458332d110c59370379a7482b4d5"
Commit 5d11ef6d authored by carlushuang's avatar carlushuang
Browse files

support dynamic buffer using memory coherence glc_slc bit from template

parent ac9e01e2
Expand all Show whitespace changes
Inline Side-by-side
Showing with 191 additions and 104 deletions
include/ck/utility/dynamic_buffer.hpp
View file @ 5d11ef6d
...@@ -52,6 +52,7 @@ struct DynamicBuffer ...@@ -52,6 +52,7 @@ struct DynamicBuffer
__host__ __device__ constexpr T& operator()(index_t i) { return p_data_[i]; } __host__ __device__ constexpr T& operator()(index_t i) { return p_data_[i]; }
template <typename X, 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 enable_if<is_same<typename scalar_type<remove_cvref_t<X>>::type,
typename scalar_type<remove_cvref_t<T>>::type>::value, typename scalar_type<remove_cvref_t<T>>::type>::value,
bool>::type = false> bool>::type = false>
...@@ -77,13 +78,16 @@ struct DynamicBuffer ...@@ -77,13 +78,16 @@ struct DynamicBuffer
if constexpr(InvalidElementUseNumericalZeroValue) if constexpr(InvalidElementUseNumericalZeroValue)
{ {
return amd_buffer_load_invalid_element_return_zero<remove_cvref_t<T>, t_per_x>( return amd_buffer_load_invalid_element_return_zero<remove_cvref_t<T>,
t_per_x,
coherence>(
p_data_, i, is_valid_element, element_space_size_); p_data_, i, is_valid_element, element_space_size_);
} }
else else
{ {
return amd_buffer_load_invalid_element_return_customized_value<remove_cvref_t<T>, return amd_buffer_load_invalid_element_return_customized_value<remove_cvref_t<T>,
t_per_x>( t_per_x,
coherence>(
p_data_, i, is_valid_element, element_space_size_, invalid_element_value_); p_data_, i, is_valid_element, element_space_size_, invalid_element_value_);
} }
} }
...@@ -144,6 +148,7 @@ struct DynamicBuffer ...@@ -144,6 +148,7 @@ struct DynamicBuffer
} }
template <typename X, 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 enable_if<is_same<typename scalar_type<remove_cvref_t<X>>::type,
typename scalar_type<remove_cvref_t<T>>::type>::value, typename scalar_type<remove_cvref_t<T>>::type>::value,
bool>::type = false> bool>::type = false>
...@@ -173,7 +178,7 @@ struct DynamicBuffer ...@@ -173,7 +178,7 @@ struct DynamicBuffer
{ {
constexpr index_t t_per_x = scalar_per_x_vector / scalar_per_t_vector; constexpr index_t t_per_x = scalar_per_x_vector / scalar_per_t_vector;
amd_buffer_store<remove_cvref_t<T>, t_per_x>( amd_buffer_store<remove_cvref_t<T>, t_per_x, coherence>(
x, p_data_, i, is_valid_element, element_space_size_); x, p_data_, i, is_valid_element, element_space_size_);
} }
else if constexpr(GetAddressSpace() == AddressSpaceEnum::Lds && else if constexpr(GetAddressSpace() == AddressSpaceEnum::Lds &&
......
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