From b9afa21694de2360e679c3a8befdad7be4d3505f Mon Sep 17 00:00:00 2001 From: Zach Bjornson Date: Sun, 14 Nov 2021 14:03:57 -0800 Subject: [PATCH] Fix aliasing violation in t-SNE Undefined behavior introduced in #4208. These parameters are restrict-qualified, but the callers pass pointers with the same address and the arrays are modified through both pointers. --- cpp/src/tsne/utils.cuh | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/cpp/src/tsne/utils.cuh b/cpp/src/tsne/utils.cuh index 8d9b68da0e..3c07a1a675 100644 --- a/cpp/src/tsne/utils.cuh +++ b/cpp/src/tsne/utils.cuh @@ -183,8 +183,8 @@ __global__ void min_max_kernel( */ template __global__ void compute_kl_div_k(const value_t* restrict Ps, - const value_t* restrict Qs, - value_t* restrict KL_divs, + const value_t* Qs, + value_t* KL_divs, const value_idx NNZ) { const auto index = (blockIdx.x * blockDim.x) + threadIdx.x; @@ -199,11 +199,8 @@ __global__ void compute_kl_div_k(const value_t* restrict Ps, * Compute KL divergence */ template -value_t compute_kl_div(value_t* restrict Ps, - value_t* restrict Qs, - value_t* restrict KL_divs, - const size_t NNZ, - cudaStream_t stream) +value_t compute_kl_div( + value_t* restrict Ps, value_t* Qs, value_t* KL_divs, const size_t NNZ, cudaStream_t stream) { value_t P_sum = thrust::reduce(rmm::exec_policy(stream), Ps, Ps + NNZ); raft::linalg::scalarMultiply(Ps, Ps, 1.0f / P_sum, NNZ, stream);