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
OpenDAS
apex
Commits
758826fc
Commit
758826fc
authored
May 11, 2020
by
Thor Johnsen
Browse files
Resolve possible race condition in stride_finite_check kernel
parent
0bfb8300
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
2 additions
and
1 deletion
+2
-1
apex/contrib/csrc/optimizers/fused_adam_cuda_kernel.cu
apex/contrib/csrc/optimizers/fused_adam_cuda_kernel.cu
+2
-1
No files found.
apex/contrib/csrc/optimizers/fused_adam_cuda_kernel.cu
View file @
758826fc
...
@@ -793,7 +793,8 @@ void fused_strided_check_finite(
...
@@ -793,7 +793,8 @@ void fused_strided_check_finite(
//Determine #threads and #blocks
//Determine #threads and #blocks
const
int
threadsPerBlock
=
512
;
const
int
threadsPerBlock
=
512
;
const
dim3
blocks
((
niter
+
threadsPerBlock
-
1
)
/
threadsPerBlock
);
//In order to avoid race condition, blocks must be 1 when clear_overflow_first flag is set.
const
dim3
blocks
(
clear_overflow_first
?
1
:
(
niter
+
threadsPerBlock
-
1
)
/
threadsPerBlock
);
AT_ASSERTM
(
at
::
cuda
::
detail
::
canUse32BitIndexMath
(
p_copy
),
"parameter tensor is too large to be indexed with int32"
);
AT_ASSERTM
(
at
::
cuda
::
detail
::
canUse32BitIndexMath
(
p_copy
),
"parameter tensor is too large to be indexed with int32"
);
cudaStream_t
stream
=
at
::
cuda
::
getCurrentCUDAStream
();
cudaStream_t
stream
=
at
::
cuda
::
getCurrentCUDAStream
();
...
...
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