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
50e573f5
Commit
50e573f5
authored
Jun 19, 2023
by
carlushuang
Browse files
fix format
parent
f74b77bc
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
18 additions
and
21 deletions
+18
-21
include/ck/utility/workgroup_barrier.hpp
include/ck/utility/workgroup_barrier.hpp
+18
-21
No files found.
include/ck/utility/workgroup_barrier.hpp
View file @
50e573f5
...
...
@@ -3,10 +3,9 @@
#include <stdint.h>
namespace
ck
{
struct
workgroup_barrier
{
__device__
workgroup_barrier
(
uint32_t
*
ptr
)
:
base_ptr
(
ptr
)
{}
struct
workgroup_barrier
{
__device__
workgroup_barrier
(
uint32_t
*
ptr
)
:
base_ptr
(
ptr
)
{}
__device__
uint32_t
ld
(
uint32_t
offset
)
{
...
...
@@ -29,48 +28,46 @@ struct workgroup_barrier {
__device__
void
wait_eq
(
uint32_t
offset
,
uint32_t
value
)
{
if
(
threadIdx
.
x
==
0
){
while
(
ld
(
offset
)
!=
value
){}
if
(
threadIdx
.
x
==
0
)
{
while
(
ld
(
offset
)
!=
value
)
{}
}
__syncthreads
();
}
__device__
void
wait_lt
(
uint32_t
offset
,
uint32_t
value
)
{
if
(
threadIdx
.
x
==
0
){
while
(
ld
(
offset
)
<
value
){}
if
(
threadIdx
.
x
==
0
)
{
while
(
ld
(
offset
)
<
value
)
{}
}
__syncthreads
();
}
__device__
void
wait_set
(
uint32_t
offset
,
uint32_t
compare
,
uint32_t
value
)
{
if
(
threadIdx
.
x
==
0
){
while
(
atomicCAS
(
base_ptr
+
offset
,
compare
,
value
)
!=
compare
){}
if
(
threadIdx
.
x
==
0
)
{
while
(
atomicCAS
(
base_ptr
+
offset
,
compare
,
value
)
!=
compare
)
{}
}
__syncthreads
();
}
// enter critical zoon, assume buffer is zero when launch kernel
__device__
void
aquire
(
uint32_t
offset
)
{
wait_set
(
offset
,
0
,
1
);
}
__device__
void
aquire
(
uint32_t
offset
)
{
wait_set
(
offset
,
0
,
1
);
}
// exit critical zoon, assume buffer is zero when launch kernel
__device__
void
release
(
uint32_t
offset
)
{
wait_set
(
offset
,
1
,
0
);
}
__device__
void
release
(
uint32_t
offset
)
{
wait_set
(
offset
,
1
,
0
);
}
__device__
void
inc
(
uint32_t
offset
)
{
__syncthreads
();
if
(
threadIdx
.
x
==
0
){
if
(
threadIdx
.
x
==
0
)
{
atomicAdd
(
base_ptr
+
offset
,
1
);
}
}
uint32_t
*
base_ptr
;
uint32_t
*
base_ptr
;
};
}
}
// namespace ck
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