Skip to content

Fix compiler warnings for 0.99.3 release #472

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

Merged
tamar-dewilde merged 4 commits into from
Aug 13, 2026
Merged

Fix compiler warnings for 0.99.3 release #472

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
e4d9bb9
Fix CPU compiler warning
tamar-dewilde Aug 12, 2026
552bb89
Fix HIP compiler warnings
tamar-dewilde Aug 13, 2026
fef1743
Apply pre-commmit fixes
tamar-dewilde Aug 13, 2026
0312f81
Address review comments
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
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