Added preconditions to avoid write-write data-race #266
Conversation
|
|
| @@ -24,6 +24,9 @@ using Complex = mindspore::utils::Complex<T>; | |||
|
|
|||
| template <typename T> | |||
| __global__ void MatrixTransposeKernel(const T *input, int elements, int row, int col, T *output) { | |||
| if (col <= 0 || row <= 0 || row != col) { | |||
There was a problem hiding this comment.
Thank you for your code committion！Could you explain why the case of non-square matrix is skipped here? And have you really encountered a situation that the col or row is less or equal than 0? If so, this should be intercepted earlier on the host side.
There was a problem hiding this comment.
Sorry about that, the constraint that the matrix is square was mistakenly added because we were looking at an old version of the kernel. I've fixed the constraint to ensure that both row/col are positive numbers.
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/matrix_transpose_impl.cu
Show resolved
Hide resolved
…to the host side API. Co-authored-by: AGroupofProbiotocs <[email protected]>
|
I think you might misunderstood my comment. What I actually mean is to move the conditional statement "if (col < 0 || row < 0 )..." into the host function "MatrixTranspose(...)" below the kernel function "MatrixTransposeKernel(...)". |
/kind bugSummary:
In issue, our tool found a write-write data-race on array output, given certain values of col,row, and elements.
We added preconditions (which may not be the minimal preconditions) to prevent the race.
mindspore-assistant