Skip to content

Release 0.99.3 #471

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
pelesh wants to merge 4 commits into
base: develop
Choose a base branch
from
Open

Release 0.99.3 #471

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
a8f8995
Update version number
pelesh Aug 12, 2026
d23596f
added Kakeru as author
shakedregev Aug 13, 2026
e7f0fe1
Update developer list in index.rst
pelesh Aug 13, 2026
d15bc2c
Fix compiler warnings for 0.99.3 release (#472)
tamar-dewilde Aug 13, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

Expand Down
1 change: 1 addition & 0 deletions README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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),
Expand Down
4 changes: 4 additions & 0 deletions docs/index.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
4 changes: 4 additions & 0 deletions examples/sysRefactor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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")
{
Expand Down
11 changes: 6 additions & 5 deletions resolve/hip/HipMemory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,6 @@
#pragma once

#include <cstddef>
#include <iostream>

#include <hip/hip_runtime.h>
Expand Down Expand Up @@ -57,7 +58,7 @@ namespace ReSolve
template <typename I, typename T>
static int allocateArrayOnDevice(T** v, I n)
{
return checkHipErrors(hipMalloc((void**) v, sizeof(T) * n));
return checkHipErrors(hipMalloc((void**) v, sizeof(T) * static_cast<std::size_t>(n)));
}

/**
Expand Down Expand Up @@ -94,7 +95,7 @@ namespace ReSolve
template <typename I, typename T>
static int setZeroArrayOnDevice(T* v, I n)
{
return checkHipErrors(hipMemset(v, 0, sizeof(T) * n));
return checkHipErrors(hipMemset(v, 0, sizeof(T) * static_cast<std::size_t>(n)));
}

/**
Expand Down Expand Up @@ -130,7 +131,7 @@ namespace ReSolve
template <typename I, typename T>
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<std::size_t>(n), hipMemcpyDeviceToHost));
}

/**
Expand All @@ -147,7 +148,7 @@ namespace ReSolve
template <typename I, typename T>
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<std::size_t>(n), hipMemcpyDeviceToDevice));
}

/**
Expand All @@ -164,7 +165,7 @@ namespace ReSolve
template <typename I, typename T>
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<std::size_t>(n), hipMemcpyHostToDevice));
}
};
} // namespace memory
Expand Down
41 changes: 20 additions & 21 deletions resolve/hip/hipKernels.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<index_type>(threadIdx.x);
index_type bsize = static_cast<index_type>(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<index_type>(blockIdx.x);
s_tmp1[t] = 0.0;
s_tmp2[t] = 0.0;
index_type nn = t;
Expand Down Expand Up @@ -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];
}
}

Expand All @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);
index_type t = static_cast<index_type>(threadIdx.x);

__shared__ real_type s_alpha[Tmaxk];
if (t < k) {
Expand Down Expand Up @@ -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<index_type>(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) {
Expand All @@ -193,11 +194,11 @@ namespace ReSolve {
real_type* result)
{

index_type idx = blockIdx.x * blockDim.x + threadIdx.x;
index_type idx = static_cast<index_type>(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<index_type>(threadIdx.x);
index_type bsize = static_cast<index_type>(blockDim.x);
real_type local_max = 0.0;
if (idx < n) {
local_max = fabs(input[idx]);
Expand Down Expand Up @@ -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];
}
}
Expand All @@ -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<index_type>(blockIdx.x*blockDim.x + threadIdx.x);
while (idx < n) {
vec_out[idx] = vec_in[perm_vector[idx]];
idx += (blockDim.x * gridDim.x);
Expand All @@ -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<index_type>(blockIdx.x*blockDim.x + threadIdx.x);
while (idx < n) {
vec_out[perm_vector[idx]] = vec_in[idx];
idx += (blockDim.x * gridDim.x);
Expand All @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);

// Check if the thread's row is within matrix bounds
if (row < n) {
Expand Down Expand Up @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);

// Check if the thread's row is within matrix bounds
if (row < n) {
Expand Down Expand Up @@ -396,7 +395,7 @@ namespace ReSolve {
real_type* result)
{
hipLaunchKernelGGL(kernels::MassIPTwoVec_kernel,
dim3(i),
dim3(static_cast<unsigned int>(i)),
dim3(1024),
0,
0,
Expand All @@ -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<unsigned int>((n + 384 - 1) / 384)),
dim3(384),
0,
0,
Expand Down Expand Up @@ -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<<<num_blocks, block_size>>>(n, a_row_ptr, a_val, d_val);
kernels::leftScale<<<static_cast<unsigned int>(num_blocks), static_cast<unsigned int>(block_size)>>>(n, a_row_ptr, a_val, d_val);
}

/**
Expand All @@ -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<<<num_blocks, block_size>>>(n, a_row_ptr, a_col_ind, a_val, d_val);
kernels::rightScale<<<static_cast<unsigned int>(num_blocks), static_cast<unsigned int>(block_size)>>>(n, a_row_ptr, a_col_ind, a_val, d_val);
}
} // namespace hip
} // namespace ReSolve
32 changes: 17 additions & 15 deletions resolve/hip/hipSketchingKernels.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
*
*
*/
#include <cstddef>

#include <hip/hip_runtime.h>
#include <hip/hip_cooperative_groups.h>
Expand Down Expand Up @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);
while (idx < n) {
real_type val = input[idx];
if (flip[idx] != 1){
Expand All @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);
while (idx < k) {
output[idx] = input[perm[idx]];
idx += blockDim.x * gridDim.x;
Expand All @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);

while (idx < n) {

Expand All @@ -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<index_type>(blockIdx.x * blockDim.x + threadIdx.x);
const index_type N = static_cast<index_type>(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<index_type>(blockIdx.y) * N;
real_type* d_Dst = d_Output + static_cast<index_type>(blockIdx.y) * N;

index_type lo = pos & (stride - 1);
index_type i0 = ((pos - lo) << 2) + lo;
Expand Down Expand Up @@ -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<index_type>(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<index_type>(threadIdx.x); pos < N; pos += static_cast<index_type>(blockDim.x)) {
s_data[pos] = d_Src[pos];
}

// Main radix-4 stages
const index_type pos = threadIdx.x;
const index_type pos = static_cast<index_type>(threadIdx.x);

for (index_type stride = N >> 2; stride > 0; stride >>= 2) {
index_type lo = pos & (stride - 1);
Expand Down Expand Up @@ -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<index_type>(threadIdx.x); pos < N / 2; pos += static_cast<index_type>(blockDim.x)) {
index_type i0 = pos << 1;
index_type i1 = i0 + 1;

Expand All @@ -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<index_type>(threadIdx.x); pos < N; pos += static_cast<index_type>(blockDim.x)) {
d_Dst[pos] = s_data[pos];
}
}
Expand Down Expand Up @@ -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<unsigned int>((1 << log2N) / (4 * THREAD_N)), static_cast<unsigned int>(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<unsigned int>(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<unsigned int>(M)), dim3(static_cast<unsigned int>(N / 4)), static_cast<std::size_t>(N) * sizeof(real_type), 0, d_Data, d_Data, log2N);
}
} // namespace hip
} // namespace ReSolve
Loading
Loading