From a8f8995543af9ef20606b937c72e9b00d3481524 Mon Sep 17 00:00:00 2001 From: pelesh Date: Wed, 12 Aug 2026 12:20:11 -0400 Subject: [PATCH 1/4] Update version number --- CMakeLists.txt | 2 +- tests/functionality/testVersion.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c366ee3c1..cee0926d2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -9,7 +9,7 @@ cmake_minimum_required(VERSION 3.22) # Adds version settings and set variable CMAKE_PROJECT_VERSION -project(ReSolve VERSION "0.99.2") +project(ReSolve VERSION "0.99.3") set(CMAKE_CXX_STANDARD 11) diff --git a/tests/functionality/testVersion.cpp b/tests/functionality/testVersion.cpp index aee29b272..2250f4b4f 100644 --- a/tests/functionality/testVersion.cpp +++ b/tests/functionality/testVersion.cpp @@ -19,7 +19,7 @@ int main() { using namespace ReSolve::colors; - std::string answer("0.99.2"); + std::string answer("0.99.3"); std::string versionstr; ReSolve::VersionGetVersionStr(versionstr); std::cout << "ReSolveVersionGetVersionStr Test: " << versionstr << std::endl From d23596f30b9d73fd0f8e08094b8afe124cc0e60b Mon Sep 17 00:00:00 2001 From: Shaked Regev Date: Thu, 13 Aug 2026 10:02:12 -0400 Subject: [PATCH 2/4] added Kakeru as author --- README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/README.md b/README.md index aa06c0850..3b72dcc50 100644 --- a/README.md +++ b/README.md @@ -134,6 +134,7 @@ Re::Solve project would not be possible without significant contributions from - Jaelyn Litzinger - Phil Roth - Cameron Rutherford +- Kakeru Ueda - Andrew Xu Development of this code was supported by the Exascale Computing Project (ECP), From e7f0fe1106836f1f7cb2c44074e11f6ad45835aa Mon Sep 17 00:00:00 2001 From: pelesh Date: Thu, 13 Aug 2026 10:26:21 -0400 Subject: [PATCH 3/4] Update developer list in index.rst --- docs/index.rst | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/docs/index.rst b/docs/index.rst index 5496c0f83..a7ff56c7f 100644 --- a/docs/index.rst +++ b/docs/index.rst @@ -50,11 +50,15 @@ from (in alphabetic order): * Maksudul Alam (ORNL) * Kaleb Brunhoeber (ORNL) * Ryan Danehy (PNNL) +* Tamar Dewilde * Adham Ibrahim (ORNL) * Nicholson Koukpaizan (ORNL) * Jaelyn Litzinger (PNNL) * Phil Roth (ORNL) * Cameron Rutherford (PNNL) +* Kakeru Ueda +* Andrew Xu + Development of this code was supported by the Exascale Computing Project (ECP), Project Number: 17-SC-20-SC, a collaborative effort of two DOE organizations From d15bc2caa1d1481cd062ff1bebf9679d66cc7165 Mon Sep 17 00:00:00 2001 From: tamar-dewilde Date: Thu, 13 Aug 2026 06:53:32 -0800 Subject: [PATCH 4/4] Fix compiler warnings for 0.99.3 release (#472) * Fix CPU compiler warning * Fix HIP compiler warnings * Apply pre-commmit fixes * Address review comments --------- Co-authored-by: tamar-dewilde --- examples/sysRefactor.cpp | 4 ++ resolve/hip/HipMemory.hpp | 11 ++--- resolve/hip/hipKernels.hip | 41 +++++++++---------- resolve/hip/hipSketchingKernels.hip | 32 ++++++++------- resolve/hip/hipVectorKernels.hip | 34 ++++++++------- .../permutation/HipPermutationKernels.hip | 4 +- resolve/hykkt/ruiz/RuizScalingKernelsHip.hip | 8 ++-- 7 files changed, 69 insertions(+), 65 deletions(-) diff --git a/examples/sysRefactor.cpp b/examples/sysRefactor.cpp index 2c850ead6..f306275cf 100644 --- a/examples/sysRefactor.cpp +++ b/examples/sysRefactor.cpp @@ -32,6 +32,10 @@ /// timing does not make an unnecessary CUDA or HIP runtime call. static void syncDevice(const std::string& hw_backend) { +#if !defined(RESOLVE_USE_CUDA) && !defined(RESOLVE_USE_HIP) + (void) hw_backend; +#endif + #ifdef RESOLVE_USE_CUDA if (hw_backend == "CUDA") { diff --git a/resolve/hip/HipMemory.hpp b/resolve/hip/HipMemory.hpp index 372da4f25..b5be9d9b7 100644 --- a/resolve/hip/HipMemory.hpp +++ b/resolve/hip/HipMemory.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include @@ -57,7 +58,7 @@ namespace ReSolve template static int allocateArrayOnDevice(T** v, I n) { - return checkHipErrors(hipMalloc((void**) v, sizeof(T) * n)); + return checkHipErrors(hipMalloc((void**) v, sizeof(T) * static_cast(n))); } /** @@ -94,7 +95,7 @@ namespace ReSolve template static int setZeroArrayOnDevice(T* v, I n) { - return checkHipErrors(hipMemset(v, 0, sizeof(T) * n)); + return checkHipErrors(hipMemset(v, 0, sizeof(T) * static_cast(n))); } /** @@ -130,7 +131,7 @@ namespace ReSolve template static int copyArrayDeviceToHost(T* dst, const T* src, I n) { - return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * n, hipMemcpyDeviceToHost)); + return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * static_cast(n), hipMemcpyDeviceToHost)); } /** @@ -147,7 +148,7 @@ namespace ReSolve template static int copyArrayDeviceToDevice(T* dst, const T* src, I n) { - return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * n, hipMemcpyDeviceToDevice)); + return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * static_cast(n), hipMemcpyDeviceToDevice)); } /** @@ -164,7 +165,7 @@ namespace ReSolve template static int copyArrayHostToDevice(T* dst, const T* src, I n) { - return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * n, hipMemcpyHostToDevice)); + return checkHipErrors(hipMemcpy(dst, src, sizeof(T) * static_cast(n), hipMemcpyHostToDevice)); } }; } // namespace memory diff --git a/resolve/hip/hipKernels.hip b/resolve/hip/hipKernels.hip index 4a9de9da1..48a958521 100644 --- a/resolve/hip/hipKernels.hip +++ b/resolve/hip/hipKernels.hip @@ -37,15 +37,15 @@ namespace ReSolve { const index_type k, const index_type N) { - index_type t = threadIdx.x; - index_type bsize = blockDim.x; + index_type t = static_cast(threadIdx.x); + index_type bsize = static_cast(blockDim.x); // assume T threads per thread block (and k reductions to be performed) volatile __shared__ real_type s_tmp1[Tv5]; volatile __shared__ real_type s_tmp2[Tv5]; // map between thread index space and the problem index space - index_type j = blockIdx.x; + index_type j = static_cast(blockIdx.x); s_tmp1[t] = 0.0; s_tmp2[t] = 0.0; index_type nn = t; @@ -113,8 +113,8 @@ namespace ReSolve { s_tmp2[t] += s_tmp2[t + 1]; } if (t == 0) { - result[blockIdx.x] = s_tmp1[0]; - result[blockIdx.x + k] = s_tmp2[0]; + result[j] = s_tmp1[0]; + result[j + k] = s_tmp2[0]; } } @@ -137,8 +137,8 @@ namespace ReSolve { real_type* y_data, const real_type* alpha) { - index_type i = blockIdx.x * blockDim.x + threadIdx.x; - index_type t = threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); + index_type t = static_cast(threadIdx.x); __shared__ real_type s_alpha[Tmaxk]; if (t < k) { @@ -170,7 +170,8 @@ namespace ReSolve { const real_type* a_val, real_type* result) { - index_type idx = blockIdx.x*blockDim.x + threadIdx.x; + (void) nnz; + index_type idx = static_cast(blockIdx.x*blockDim.x + threadIdx.x); while (idx < n) { real_type sum = 0.0; for (index_type i = a_ia[idx]; i < a_ia[idx+1]; ++i) { @@ -193,11 +194,11 @@ namespace ReSolve { real_type* result) { - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); volatile __shared__ real_type s_max[1024]; - index_type t = threadIdx.x; - index_type bsize = blockDim.x; + index_type t = static_cast(threadIdx.x); + index_type bsize = static_cast(blockDim.x); real_type local_max = 0.0; if (idx < n) { local_max = fabs(input[idx]); @@ -252,8 +253,6 @@ namespace ReSolve { } if (t == 0) { - index_type bid = blockIdx.x; - index_type gid = gridDim.x; result[blockIdx.x] = s_max[0]; } } @@ -273,7 +272,7 @@ namespace ReSolve { real_type* vec_out) { //one thread per vector entry, pass through rows - index_type idx = blockIdx.x*blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x*blockDim.x + threadIdx.x); while (idx < n) { vec_out[idx] = vec_in[perm_vector[idx]]; idx += (blockDim.x * gridDim.x); @@ -294,7 +293,7 @@ namespace ReSolve { real_type* vec_out) { //one thread per vector entry, pass through rows - index_type idx = blockIdx.x*blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x*blockDim.x + threadIdx.x); while (idx < n) { vec_out[perm_vector[idx]] = vec_in[idx]; idx += (blockDim.x * gridDim.x); @@ -317,7 +316,7 @@ namespace ReSolve { const real_type* d_val) { // Get row index from thread and block indices - index_type row = blockIdx.x * blockDim.x + threadIdx.x; + index_type row = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the thread's row is within matrix bounds if (row < n) { @@ -353,7 +352,7 @@ namespace ReSolve { const real_type* d_val) { // Get row index from thread and block indices - index_type row = blockIdx.x * blockDim.x + threadIdx.x; + index_type row = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the thread's row is within matrix bounds if (row < n) { @@ -396,7 +395,7 @@ namespace ReSolve { real_type* result) { hipLaunchKernelGGL(kernels::MassIPTwoVec_kernel, - dim3(i), + dim3(static_cast(i)), dim3(1024), 0, 0, @@ -420,7 +419,7 @@ namespace ReSolve { void axpy_multi(index_type n, index_type i, real_type* x, real_type* y, real_type* alpha) { hipLaunchKernelGGL(kernels::axpyMulti3_kernel, - dim3((n + 384 - 1) / 384), + dim3(static_cast((n + 384 - 1) / 384)), dim3(384), 0, 0, @@ -539,7 +538,7 @@ namespace ReSolve { const int block_size = 1; int num_blocks = (n + block_size - 1) / block_size; // Launch the kernel - kernels::leftScale<<>>(n, a_row_ptr, a_val, d_val); + kernels::leftScale<<(num_blocks), static_cast(block_size)>>>(n, a_row_ptr, a_val, d_val); } /** @@ -563,7 +562,7 @@ namespace ReSolve { const int block_size = 256; int num_blocks = (n + block_size - 1) / block_size; // Launch the kernel - kernels::rightScale<<>>(n, a_row_ptr, a_col_ind, a_val, d_val); + kernels::rightScale<<(num_blocks), static_cast(block_size)>>>(n, a_row_ptr, a_col_ind, a_val, d_val); } } // namespace hip } // namespace ReSolve diff --git a/resolve/hip/hipSketchingKernels.hip b/resolve/hip/hipSketchingKernels.hip index 4fabceaf1..de2c10af3 100644 --- a/resolve/hip/hipSketchingKernels.hip +++ b/resolve/hip/hipSketchingKernels.hip @@ -5,6 +5,7 @@ * * */ +#include #include #include @@ -34,7 +35,8 @@ namespace ReSolve const real_type* input, real_type* output) { - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + (void) k; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); while (idx < n) { real_type val = input[idx]; if (flip[idx] != 1){ @@ -58,7 +60,7 @@ namespace ReSolve const real_type* input, real_type* output) { - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); while (idx < k) { output[idx] = input[perm[idx]]; idx += blockDim.x * gridDim.x; @@ -78,7 +80,7 @@ namespace ReSolve const real_type* x, real_type* y) { - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); while (idx < n) { @@ -101,11 +103,11 @@ namespace ReSolve */ __global__ void fwtBatch2Kernel(real_type* d_Output, real_type* d_Input, index_type stride) { - const index_type pos = blockIdx.x * blockDim.x + threadIdx.x; - const index_type N = blockDim.x * gridDim.x * 4; + const index_type pos = static_cast(blockIdx.x * blockDim.x + threadIdx.x); + const index_type N = static_cast(blockDim.x * gridDim.x * 4); - real_type* d_Src = d_Input + blockIdx.y * N; - real_type* d_Dst = d_Output + blockIdx.y * N; + real_type* d_Src = d_Input + static_cast(blockIdx.y) * N; + real_type* d_Dst = d_Output + static_cast(blockIdx.y) * N; index_type lo = pos & (stride - 1); index_type i0 = ((pos - lo) << 2) + lo; @@ -150,19 +152,19 @@ namespace ReSolve cooperative_groups::thread_block cta = cooperative_groups::this_thread_block(); const index_type N = 1 << log2N; - const index_type base = blockIdx.x << log2N; + const index_type base = static_cast(blockIdx.x) << log2N; //(2 ** 11) * 4 bytes == 8KB -- maximum s_data[] size for G80 extern __shared__ real_type s_data[]; real_type* d_Src = d_Input + base; real_type* d_Dst = d_Output + base; - for (index_type pos = threadIdx.x; pos < N; pos += blockDim.x) { + for (index_type pos = static_cast(threadIdx.x); pos < N; pos += static_cast(blockDim.x)) { s_data[pos] = d_Src[pos]; } // Main radix-4 stages - const index_type pos = threadIdx.x; + const index_type pos = static_cast(threadIdx.x); for (index_type stride = N >> 2; stride > 0; stride >>= 2) { index_type lo = pos & (stride - 1); @@ -197,7 +199,7 @@ namespace ReSolve cooperative_groups::sync(cta); - for (index_type pos = threadIdx.x; pos < N / 2; pos += blockDim.x) { + for (index_type pos = static_cast(threadIdx.x); pos < N / 2; pos += static_cast(blockDim.x)) { index_type i0 = pos << 1; index_type i1 = i0 + 1; @@ -210,7 +212,7 @@ namespace ReSolve cooperative_groups::sync(cta); - for (index_type pos = threadIdx.x; pos < N; pos += blockDim.x) { + for (index_type pos = static_cast(threadIdx.x); pos < N; pos += static_cast(blockDim.x)) { d_Dst[pos] = s_data[pos]; } } @@ -303,13 +305,13 @@ namespace ReSolve const index_type ELEMENTARY_LOG2SIZE = 11; const index_type THREAD_N = 1024; index_type N = 1 << log2N; - dim3 grid((1 << log2N) / (4 * THREAD_N), M, 1); + dim3 grid(static_cast((1 << log2N) / (4 * THREAD_N)), static_cast(M), 1); for (; log2N > ELEMENTARY_LOG2SIZE; log2N -= 2, N >>= 2, M <<= 2) { - hipLaunchKernelGGL(kernels::fwtBatch2Kernel, grid, dim3(THREAD_N), 0, 0, d_Data, d_Data, N / 4); + hipLaunchKernelGGL(kernels::fwtBatch2Kernel, grid, dim3(static_cast(THREAD_N)), 0, 0, d_Data, d_Data, N / 4); } - hipLaunchKernelGGL(kernels::fwtBatch1Kernel, dim3(M), dim3(N / 4), N * sizeof(real_type), 0, d_Data, d_Data, log2N); + hipLaunchKernelGGL(kernels::fwtBatch1Kernel, dim3(static_cast(M)), dim3(static_cast(N / 4)), static_cast(N) * sizeof(real_type), 0, d_Data, d_Data, log2N); } } // namespace hip } // namespace ReSolve diff --git a/resolve/hip/hipVectorKernels.hip b/resolve/hip/hipVectorKernels.hip index 8242e314d..f2626be6f 100644 --- a/resolve/hip/hipVectorKernels.hip +++ b/resolve/hip/hipVectorKernels.hip @@ -20,7 +20,7 @@ namespace ReSolve { __global__ void set_array_to_const(index_type n, real_type val, real_type* arr) { - index_type i = blockIdx.x * blockDim.x + threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); while (i < n) { arr[i] = val; @@ -40,7 +40,7 @@ namespace ReSolve { */ __global__ void addConst(index_type n, real_type val, real_type* arr) { - index_type i = blockIdx.x * blockDim.x + threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); if(i < n) { arr[i] += val; @@ -61,7 +61,7 @@ namespace ReSolve { real_type* vec) { // Get the index of the element to be processed - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the index is within bounds if (idx < n) { @@ -84,7 +84,7 @@ namespace ReSolve { real_type* vec) { // Get the index of the element to be processed - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the index is within bounds if (idx < n) @@ -110,7 +110,7 @@ namespace ReSolve { real_type* out) { // Get the index of the element to be processed - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the index is within bounds if (idx < n) @@ -134,7 +134,7 @@ namespace ReSolve { real_type* out) { // Get the index of the element to be processed - index_type idx = blockIdx.x * blockDim.x + threadIdx.x; + index_type idx = static_cast(blockIdx.x * blockDim.x + threadIdx.x); // Check if the index is within bounds if (idx < n) @@ -145,21 +145,19 @@ namespace ReSolve { } } // namespace kernels - constexpr index_type block_size = 256; + constexpr size_t block_size = 256; void setArrayConst(index_type n, real_type c, real_type* v) { - index_type num_blocks; - index_type block_size = 512; - num_blocks = (n + block_size - 1) / block_size; + const size_t block_size = 512; + const size_t num_blocks = (n + block_size - 1) / block_size; hipLaunchKernelGGL(kernels::set_array_to_const, dim3(num_blocks), dim3(block_size), 0, 0, n, c, v); } void addConst(index_type n, real_type val, real_type* arr) { - index_type num_blocks; - index_type block_size = 512; - num_blocks = (n + block_size - 1) / block_size; + const size_t block_size = 512; + const size_t num_blocks = (n + block_size - 1) / block_size; hipLaunchKernelGGL(kernels::addConst, dim3(num_blocks), dim3(block_size), 0, 0, n, val, arr); } @@ -177,8 +175,8 @@ namespace ReSolve { real_type* vec) { // Define block size and number of blocks - const int block_size = 256; - int num_blocks = (n + block_size - 1) / block_size; + const size_t block_size = 256; + const size_t num_blocks = (n + block_size - 1) / block_size; // Launch the kernel kernels::scale<<>>(n, diag, vec); } @@ -196,7 +194,7 @@ namespace ReSolve { const real_type* diag, real_type* vec) { - int num_blocks = (n + block_size - 1) / block_size; + const size_t num_blocks = (n + block_size - 1) / block_size; // Launch the kernel kernels::diagSolve<<>>(n, diag, vec); } @@ -216,7 +214,7 @@ namespace ReSolve { const real_type* y, real_type* out) { - int num_blocks = (n + block_size - 1) / block_size; + const size_t num_blocks = (n + block_size - 1) / block_size; // Launch the kernel kernels::max<<>>(n, x, y, out); } @@ -234,7 +232,7 @@ namespace ReSolve { const real_type* in, real_type* out) { - int num_blocks = (n + block_size - 1) / block_size; + const size_t num_blocks = (n + block_size - 1) / block_size; // Launch the kernel kernels::abs<<>>(n, in, out); } diff --git a/resolve/hykkt/permutation/HipPermutationKernels.hip b/resolve/hykkt/permutation/HipPermutationKernels.hip index 939f62f83..9e7aed9ed 100644 --- a/resolve/hykkt/permutation/HipPermutationKernels.hip +++ b/resolve/hykkt/permutation/HipPermutationKernels.hip @@ -22,7 +22,7 @@ namespace ReSolve */ __global__ void mapIdxKernel(index_type n, const index_type* perm, const real_type* old_val, real_type* new_val) { - index_type i = blockIdx.x * blockDim.x + threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); if (i < n) { new_val[i] = old_val[perm[i]]; @@ -47,7 +47,7 @@ namespace ReSolve // Launch the CUDA kernel index_type blockSize = 256; index_type numBlocks = (n + blockSize - 1) / blockSize; - mapIdxKernel<<>>(n, perm, old_val, new_val); + mapIdxKernel<<(numBlocks), static_cast(blockSize)>>>(n, perm, old_val, new_val); } } // namespace hykkt } // namespace ReSolve diff --git a/resolve/hykkt/ruiz/RuizScalingKernelsHip.hip b/resolve/hykkt/ruiz/RuizScalingKernelsHip.hip index 177e4a530..ca2d9afb0 100644 --- a/resolve/hykkt/ruiz/RuizScalingKernelsHip.hip +++ b/resolve/hykkt/ruiz/RuizScalingKernelsHip.hip @@ -36,7 +36,7 @@ namespace ReSolve { real_type max_l = 0; real_type max_u = 0; - index_type i = blockIdx.x * blockDim.x + threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); index_type j; real_type entry; if (i < n_hes) @@ -115,7 +115,7 @@ namespace ReSolve real_type* aggregate_scaling_vector, const real_type* scaling_vector) { - index_type i = blockIdx.x * blockDim.x + threadIdx.x; + index_type i = static_cast(blockIdx.x * blockDim.x + threadIdx.x); if (i < n_hes) { for (index_type j = hes_i[i]; j < hes_i[i + 1]; j++) @@ -169,7 +169,7 @@ namespace ReSolve int num_blocks; int block_size = 256; num_blocks = (n_total + block_size - 1) / block_size; - kernelsHIP::adaptRowMax<<>>(n_hes, + kernelsHIP::adaptRowMax<<(num_blocks), static_cast(block_size)>>>(n_hes, n_total, hes->getRowData(memory::DEVICE), hes->getValues(memory::DEVICE), @@ -214,7 +214,7 @@ namespace ReSolve { int block_size = 256; int num_blocks = (n_total + block_size - 1) / block_size; - kernelsHIP::adaptDiagScale<<>>(n_hes, + kernelsHIP::adaptDiagScale<<(num_blocks), static_cast(block_size)>>>(n_hes, n_total, hes->getRowData(memory::DEVICE), hes->getColData(memory::DEVICE),