Add support of cuda kernels as pullback functions #1114
Conversation
Codecov ReportAttention: Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## master #1114 +/- ##
==========================================
+ Coverage 94.36% 94.38% +0.01%
==========================================
Files 50 50
Lines 8365 8744 +379
==========================================
+ Hits 7894 8253 +359
- Misses 471 491 +20
... and 26 files with indirect coverage changes
|
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
vgvassilev
force-pushed
the
cuda-pullback-host-functions
branch
from
22399c4
to
ed7194c
Compare
kchristin22
force-pushed
the
cuda-pullback-host-functions
branch
from
ed7194c
to
96ad6a1
Compare
kchristin22
force-pushed
the
cuda-pullback-host-functions
branch
from
96ad6a1
to
e5fa81d
Compare
test/CUDA/GradientKernels.cu
Outdated
| //CHECK-NEXT: cudaMemcpyKind _r1 = static_cast<cudaMemcpyKind>(0U); | ||
| //CHECK-NEXT: clad::custom_derivatives::cudaMemcpy_pullback(out_host, out, 10 * sizeof(double), cudaMemcpyDeviceToHost, _d_out_host, _d_out, &_r0, &_r1); | ||
| //CHECK-NEXT: } | ||
| //CHECK-NEXT: kernel_call_pullback<<<1, 10>>>(out, in, _d_out, _d_in); |
There was a problem hiding this comment.
Should we call cudaDeviceSynchronize before cudaFree calls?
There was a problem hiding this comment.
These should be executed in order so no need
Edit: meaning the previous GPU ops will finish before free is executed as they are all executed in the default stream. Maybe it would be worth it to add them in case the user uses another stream for the kernel, as they should also include cudaDeviceSynchronize() in the original kernel. What do you think?
There was a problem hiding this comment.
meaning the previous GPU ops will finish before free is executed as they are all executed in the default stream.
What if there are no cudaFree calls? Should we call cudaDeviceSynchronize in that case?
Maybe it would be worth it to add them in case the user uses another stream for the kernel, as they should also include cudaDeviceSynchronize() in the original kernel. What do you think?
I believe we should not add additional cudaDeviceSynchronize calls because of multiple streams currently. We can add support for multiple stream by adding extra synchronization calls if the need arises.
There was a problem hiding this comment.
They are supported for derivative kernel execution though. I can remove the cudaDeviceSynchronize calls from cudaMemcpy_pullback, but I think it doesn't hurt to keep them, as they will be added in the future either way.
|
clang-tidy review says "All clean, LGTM! 👍" |
…y_pullback to mimic its blocking behavior
|
clang-tidy review says "All clean, LGTM! 👍" |
|
clang-tidy review says "All clean, LGTM! 👍" |
No description provided.