Fix CUDA gradient script #806
Conversation
Codecov ReportAll modified and coverable lines are covered by tests ✅
Additional details and impacted files@@ Coverage Diff @@
## master #806 +/- ##
=======================================
Coverage 94.89% 94.89%
=======================================
Files 49 49
Lines 7478 7478
=======================================
Hits 7096 7096
Misses 382 382 |
|
It would be also safer to add checks for CUDA errors when calling CUDA functions, but the final comparison of the GPU and CPU results will hint on a need for debugging, so I believe there's no need to "clutter" the code further. |
| @@ -15,6 +15,7 @@ | |||
| // XFAIL: clang-15 | |||
There was a problem hiding this comment.
Can we understand why it fails for clang-15?
There was a problem hiding this comment.
My first guess would be that the local CUDA version where this was tested and failed was not compatible with clang-15. I will have a look.
There was a problem hiding this comment.
Unfortunately, I haven't found a CUDA version old enough (prior to 11.5) and compatible with my system (Ubuntu 22.04) to test this with clang-15, but upon examining another issue I came across a similar failure of clang-15 and clang-14 (didn't test with clang-16). Relevant comment here.
There was a problem hiding this comment.
This looks good to me. We probably should open a new issue and a separate PR to continue the XFAIL: clang-15 discussion there.
There was a problem hiding this comment.
I can open one if you'd like.
I was also thinking that if the CI was expanded to include CUDA's set up, we could check the build errors in the PR (harder for development but for older versions of clang it would be beneficial for developers with the latest OS versions). This could be a separate issue.
test/CUDA/GradientCuda.cu
Outdated
| @@ -109,14 +110,24 @@ int main(void) { | |||
| cudaMemcpy(d_x, x, N * sizeof(double), cudaMemcpyHostToDevice); | |||
| cudaMalloc(&d_p, N * sizeof(double)); | |||
| cudaMemcpy(d_p, p, N * sizeof(double), cudaMemcpyHostToDevice); | |||
| double *result, *d_result; | |||
| std::vector<double> result(N, 0); | |||
There was a problem hiding this comment.
We probably do not need a std::vector here. Can we use a fixed size array or an std::array?
| @@ -15,6 +15,7 @@ | |||
| // XFAIL: clang-15 | |||
There was a problem hiding this comment.
This looks good to me. We probably should open a new issue and a separate PR to continue the XFAIL: clang-15 discussion there.
Fixes #744.
As described here, the gradient computation of the function of interest was moved inside the kernel. The results returned from the GPU were compared with the ones derived from the CPU execution to establish the correctness of the implementation.