jeffdaily commented on PR #1399: URL: https://github.com/apache/mahout/pull/1399#issuecomment-4682322151
Thanks for the review, and for running the CUDA path on the 2080 Ti. All four addressed in 0b5042e: - Drop now binds the owning device before `hipFree` (the merge-blocker). - The `hipMemoryType` check reads `hipPointerAttribute_t` and compares against `hipMemoryTypeDevice` instead of the hardcoded `2`. - `build.rs` fails loudly when `QDP_USE_HIP` and the `hip` feature disagree. - The stream is now non-blocking (`hipStreamCreateWithFlags(.., hipStreamNonBlocking)`) to match cudarc. On #4: making the stream non-blocking surfaced a pre-existing latent race in the batch-f32 amplitude path -- it read the norm back on the default stream without syncing the caller's stream, masked until now by the blocking stream. Fixed with the same stream-sync the other batch paths use; it's shared (non-HIP) code, so it applies to the CUDA path too. On byte-for-byte: that's overstated, now corrected in the PR description. `metrics.rs` swapped the driver `cuMemcpyDtoH_v2` for the runtime `cudaMemcpy`, the kernel `>>5` became `/warpSize`, and the #4 sync adds one more -- all SASS-changing but behavior-identical. The CUDA path is behavior-preserving, not literally byte/SASS-identical. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
