Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
790467d6
Commit
790467d6
authored
May 25, 2023
by
Qianfeng Zhang
Committed by
carlushuang
May 26, 2023
Browse files
Add template parameter to make_dynamic_buffer to support amd_buffer coherence setting
parent
4d1a922e
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
3 deletions
+8
-3
include/ck/utility/dynamic_buffer.hpp
include/ck/utility/dynamic_buffer.hpp
+8
-3
No files found.
include/ck/utility/dynamic_buffer.hpp
View file @
790467d6
...
@@ -380,14 +380,19 @@ struct DynamicBuffer
...
@@ -380,14 +380,19 @@ struct DynamicBuffer
__host__
__device__
static
constexpr
bool
IsDynamicBuffer
()
{
return
true
;
}
__host__
__device__
static
constexpr
bool
IsDynamicBuffer
()
{
return
true
;
}
};
};
template
<
AddressSpaceEnum
BufferAddressSpace
,
typename
T
,
typename
ElementSpaceSize
>
template
<
AddressSpaceEnum
BufferAddressSpace
,
amd_buffer_coherence_bits
coherence
=
amd_buffer_coherence_bits
::
default_coherence
,
typename
T
,
typename
ElementSpaceSize
>
__host__
__device__
constexpr
auto
make_dynamic_buffer
(
T
*
p
,
ElementSpaceSize
element_space_size
)
__host__
__device__
constexpr
auto
make_dynamic_buffer
(
T
*
p
,
ElementSpaceSize
element_space_size
)
{
{
return
DynamicBuffer
<
BufferAddressSpace
,
T
,
ElementSpaceSize
,
true
>
{
p
,
element_space_size
};
return
DynamicBuffer
<
BufferAddressSpace
,
T
,
ElementSpaceSize
,
true
,
coherence
>
{
p
,
element_space_size
};
}
}
template
<
template
<
AddressSpaceEnum
BufferAddressSpace
,
AddressSpaceEnum
BufferAddressSpace
,
amd_buffer_coherence_bits
coherence
=
amd_buffer_coherence_bits
::
default_coherence
,
typename
T
,
typename
T
,
typename
ElementSpaceSize
,
typename
ElementSpaceSize
,
typename
X
,
typename
X
,
...
@@ -395,7 +400,7 @@ template <
...
@@ -395,7 +400,7 @@ template <
__host__
__device__
constexpr
auto
__host__
__device__
constexpr
auto
make_dynamic_buffer
(
T
*
p
,
ElementSpaceSize
element_space_size
,
X
invalid_element_value
)
make_dynamic_buffer
(
T
*
p
,
ElementSpaceSize
element_space_size
,
X
invalid_element_value
)
{
{
return
DynamicBuffer
<
BufferAddressSpace
,
T
,
ElementSpaceSize
,
false
>
{
return
DynamicBuffer
<
BufferAddressSpace
,
T
,
ElementSpaceSize
,
false
,
coherence
>
{
p
,
element_space_size
,
invalid_element_value
};
p
,
element_space_size
,
invalid_element_value
};
}
}
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment