diff --git a/src/LinAlg/VectorCudaKernels.cu b/src/LinAlg/VectorCudaKernels.cu index 09948dff4..405dfff78 100644 --- a/src/LinAlg/VectorCudaKernels.cu +++ b/src/LinAlg/VectorCudaKernels.cu @@ -707,9 +707,9 @@ __global__ void relax_bounds_cu(int n, /** @brief set d_ptr[i] = 1 if d1[i] == d2[i], otherwirse 0 */ __global__ void set_if_match_cu(int n, - double* d_ptr, - double* d1, - double* d2) + int* d_ptr, + const double* d1, + const double* d2) { const int num_threads = blockDim.x * gridDim.x; const int tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -1358,7 +1358,7 @@ void copyToStartingAt_w_pattern_kernel(int n_src, dd); } -int num_match_local_kernel(int n, double* d1, const double* d2) +int num_match_local_kernel(int n, const double* d1, const double* d2) { int num_blocks = (n+block_size-1)/block_size; diff --git a/src/LinAlg/VectorCudaKernels.hpp b/src/LinAlg/VectorCudaKernels.hpp index 1afe55b82..3d8cf6cba 100644 --- a/src/LinAlg/VectorCudaKernels.hpp +++ b/src/LinAlg/VectorCudaKernels.hpp @@ -297,7 +297,7 @@ void copyToStartingAt_w_pattern_kernel(int n_src, const double* dd); /// @brief return the numbers of identical elements between two vectors -int num_match_local_kernel(int n, double* d1, const double* d2); +int num_match_local_kernel(int n, const double* d1, const double* d2); /** @brief process variable bounds */ void process_bounds_local_kernel(int n_local, diff --git a/src/LinAlg/VectorHipKernels.cpp b/src/LinAlg/VectorHipKernels.cpp index 85acb17b0..bcea53a66 100644 --- a/src/LinAlg/VectorHipKernels.cpp +++ b/src/LinAlg/VectorHipKernels.cpp @@ -702,9 +702,9 @@ __global__ void relax_bounds_hip(int n, /** @brief set d_ptr[i] = 1 if d1[i] == d2[i], otherwirse 0 */ __global__ void set_if_match_hip(int n, - double* d_ptr, - double* d1, - double* d2) + int* d_ptr, + const double* d1, + const double* d2) { const int num_threads = blockDim.x * gridDim.x; const int tid = blockIdx.x * blockDim.x + threadIdx.x; @@ -1353,7 +1353,7 @@ void copyToStartingAt_w_pattern_kernel(int n_src, dd); } -int num_match_local_kernel(int n, double* d1, const double* d2) +int num_match_local_kernel(int n, const double* d1, const double* d2) { int num_blocks = (n+block_size-1)/block_size; @@ -1362,7 +1362,7 @@ int num_match_local_kernel(int n, double* d1, const double* d2) int* d_ptr = thrust::raw_pointer_cast(dv_ptr); // set d_ptr[i] = 1 if d1[i] == d2[i], otherwirse 0 - set_if_match_cu<<>>(n, d_ptr, d1, d2); + set_if_match_hip<<>>(n, d_ptr, d1, d2); int rval = thrust::reduce(thrust::device, d_ptr, d_ptr+n, 0, thrust::plus()); diff --git a/src/LinAlg/VectorHipKernels.hpp b/src/LinAlg/VectorHipKernels.hpp index 5d851e2de..ece799f18 100644 --- a/src/LinAlg/VectorHipKernels.hpp +++ b/src/LinAlg/VectorHipKernels.hpp @@ -297,7 +297,7 @@ void copyToStartingAt_w_pattern_kernel(int n_src, const double* dd); /// @brief return the numbers of identical elements between two vectors -int num_match_local_kernel(int n, double* d1, const double* d2); +int num_match_local_kernel(int n, const double* d1, const double* d2); /** @brief process variable bounds */ void process_bounds_local_kernel(int n_local, diff --git a/src/LinAlg/hiopVectorCuda.cpp b/src/LinAlg/hiopVectorCuda.cpp index ea0d14b0c..1bf877992 100644 --- a/src/LinAlg/hiopVectorCuda.cpp +++ b/src/LinAlg/hiopVectorCuda.cpp @@ -1089,8 +1089,8 @@ bool hiopVectorCuda::is_equal(const hiopVector& vec) const size_type hiopVectorCuda::num_match(const hiopVector& vec) const { - double* dd = data_; - double* vd = vec.local_data(); + const double* dd = data_; + const double* vd = vec.local_data_const(); int sum_match = sum_match = hiop::cuda::num_match_local_kernel(n_local_, dd, vd); #ifdef HIOP_USE_MPI @@ -1130,6 +1130,8 @@ bool hiopVectorCuda::process_bounds_local(const hiopVector& xu, n_bnds_lu, n_fixed_vars, fixed_var_tol); + + return true; } void hiopVectorCuda::relax_bounds_vec(hiopVector& xu, diff --git a/src/LinAlg/hiopVectorHip.cpp b/src/LinAlg/hiopVectorHip.cpp index 4f6a144de..e2a173468 100644 --- a/src/LinAlg/hiopVectorHip.cpp +++ b/src/LinAlg/hiopVectorHip.cpp @@ -1093,8 +1093,8 @@ bool hiopVectorHip::is_equal(const hiopVector& vec) const size_type hiopVectorHip::num_match(const hiopVector& vec) const { - double* dd = data_; - double* vd = vec.local_data(); + const double* dd = data_; + const double* vd = vec.local_data_const(); int sum_match = sum_match = hiop::cuda::num_match_local_kernel(n_local_, dd, vd); #ifdef HIOP_USE_MPI @@ -1134,6 +1134,8 @@ bool hiopVectorHip::process_bounds_local(const hiopVector& xu, n_bnds_lu, n_fixed_vars, fixed_var_tol); + + return true; } void hiopVectorHip::relax_bounds_vec(hiopVector& xu, diff --git a/src/LinAlg/hiopVectorRajaImpl.hpp b/src/LinAlg/hiopVectorRajaImpl.hpp index ff3091162..6e8b552cc 100644 --- a/src/LinAlg/hiopVectorRajaImpl.hpp +++ b/src/LinAlg/hiopVectorRajaImpl.hpp @@ -2250,7 +2250,7 @@ size_type hiopVectorRaja::num_match(const hiopVector& vec) const sum += 1; } }); - int all_equal = (sum.get() == 0); + int all_equal = sum.get(); #ifdef HIOP_USE_MPI int all_equalG; diff --git a/tests/LinAlg/vectorTests.hpp b/tests/LinAlg/vectorTests.hpp index 1d81711b5..78d2f820e 100644 --- a/tests/LinAlg/vectorTests.hpp +++ b/tests/LinAlg/vectorTests.hpp @@ -2058,7 +2058,6 @@ class VectorTests : public TestBase { const local_ordinal_type Nx = x.get_size(); int fail = 0; - int n_match = 0; x.setToConstant(one); y.setToConstant(one);