AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Remove comparisons of signed to unsigned values

Browse Source
This commit is contained in:
Yiltan Temucin
2024-12-12 10:21:08 -06:00
parent 1b5a6e74aa
commit fa0858833e
22 changed files with 74 additions and 74 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/constants.hpp
+1 -1
View File
@@ -55,7 +55,7 @@ inline const unsigned MAX_WG_SIZE{1024};
*
* @note Wavefront size on most systems is either 32 or 64.
*/
inline const unsigned WF_SIZE{64};
inline const int WF_SIZE{64};
} // namespace rocshmem
src/ipc/backend_ipc.cpp
+8 -8
View File
@@ -137,7 +137,7 @@ IPCBackend::~IPCBackend() {
void IPCBackend::setup_ctxs() {
CHECK_HIP(hipMalloc(&ctx_array, sizeof(IPCContext) * maximum_num_contexts_));
for (int i = 0; i < maximum_num_contexts_; i++) {
for (size_t i = 0; i < maximum_num_contexts_; i++) {
new (&ctx_array[i]) IPCContext(this);
ctx_free_list.get()->push_back(ctx_array + i);
}
@@ -367,7 +367,7 @@ void IPCBackend::init_wrk_sync_buffer() {
* For all local processing elements, initialize the device-side array
* with the IPC work/sync buffer addresses.
*/
for (size_t i = 0; i < num_pes; i++) {
for (int i = 0; i < num_pes; i++) {
if (i != my_pe) {
CHECK_HIP(hipIpcOpenMemHandle(
reinterpret_cast<void**>(&Wrk_Sync_buffer_bases_[i]),
@@ -380,7 +380,7 @@ void IPCBackend::init_wrk_sync_buffer() {
}
void IPCBackend::cleanup_wrk_sync_buffer() {
for (size_t i = 0; i < num_pes; i++) {
for (int i = 0; i < num_pes; i++) {
if (i != my_pe) {
CHECK_HIP(hipIpcCloseMemHandle(Wrk_Sync_buffer_bases_[i]));
}
@@ -444,7 +444,7 @@ void IPCBackend::teams_init() {
/* Accommodating for largest possible data type for pWrk */
pWrk_pool = reinterpret_cast<void *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE
* max_num_teams;
@@ -466,16 +466,16 @@ void IPCBackend::teams_init() {
alltoall_pSync = reinterpret_cast<long *>(
&alltoall_pSync_pool[team_i * ROCSHMEM_ALLTOALL_SYNC_SIZE]);
for (int i = 0; i < ROCSHMEM_BARRIER_SYNC_SIZE; i++) {
for (size_t i = 0; i < ROCSHMEM_BARRIER_SYNC_SIZE; i++) {
barrier_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROCSHMEM_REDUCE_SYNC_SIZE; i++) {
for (size_t i = 0; i < ROCSHMEM_REDUCE_SYNC_SIZE; i++) {
reduce_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROCSHMEM_BCAST_SYNC_SIZE; i++) {
for (size_t i = 0; i < ROCSHMEM_BCAST_SYNC_SIZE; i++) {
bcast_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROCSHMEM_ALLTOALL_SYNC_SIZE; i++) {
for (size_t i = 0; i < ROCSHMEM_ALLTOALL_SYNC_SIZE; i++) {
alltoall_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
}
src/ipc/context_ipc_device_coll.cpp
+3 -3
View File
@@ -38,14 +38,14 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
for (size_t i = 1; i < n_pes; i++) {
for (int i = 1; i < n_pes; i++) {
wait_until(&pSync[i], ROCSHMEM_CMP_EQ, flag_val);
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
threadfence_system();
// Announce to other PEs that all have reached
for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
for (int i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
internal_putmem(&pSync[0], &flag_val, sizeof(*pSync), j);
#if defined(__gfx90a__)
__threadfence_system();
@@ -73,7 +73,7 @@ __device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
for (int i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
internal_putmem(&pSync[0], &flag_val, sizeof(*pSync), j);
}
} else {
src/ipc/context_ipc_tmpl_device.hpp
+5 -5
View File
@@ -156,7 +156,7 @@ __device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) {
template <typename T, ROCSHMEM_OP Op>
__device__ void compute_reduce(T *src, T *dst, int size, int wg_id,
int wg_size) {
for (size_t i = wg_id; i < size; i += wg_size) {
for (int i = wg_id; i < size; i += wg_size) {
OpWrap<Op>::Calc(src, dst, i);
}
__syncthreads();
@@ -299,12 +299,12 @@ __device__ void IPCContext::internal_ring_allreduce(
int wg_size = get_flat_block_size();
int wg_id = get_flat_block_id();
for (size_t i = wg_id; i < nelems; i += wg_size) {
for (int i = wg_id; i < nelems; i += wg_size) {
dst[i] = src[i];
}
__syncthreads();
for (size_t seg = 0; seg < n_seg; seg++) {
for (int seg = 0; seg < n_seg; seg++) {
off_seg = seg * seg_size;
// Loop 2 in the algorithm above
for (int iter = 0; iter < PE_size - 1; iter++) {
@@ -331,7 +331,7 @@ __device__ void IPCContext::internal_ring_allreduce(
}
// Loop 2 in the example above
for (size_t iter = PE_size - 1; iter < 2 * PE_size - 2; iter++) {
for (int iter = PE_size - 1; iter < 2 * PE_size - 2; iter++) {
off_send = (((my_pe_in_team + 1 - iter + 2 * PE_size) % PE_size) * chunk_size);
putmem_nbi_wg(reinterpret_cast<void *>(&dst[off_send + off_seg]),
reinterpret_cast<void *>(&dst[off_send + off_seg]),
@@ -351,7 +351,7 @@ __device__ void IPCContext::internal_ring_allreduce(
}
__syncthreads();
for (size_t i = wg_id; i < 2 * num_pes - 2; i += wg_size) {
for (int i = wg_id; i < 2 * num_pes - 2; i += wg_size) {
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
__syncthreads();
src/ipc_policy.cpp
+2 -2
View File
@@ -86,7 +86,7 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
* For all local processing elements, initialize the device-side array
* with the IPC symmetric heap base addresses.
*/
for (size_t i = 0; i < shm_size; i++) {
for (int i = 0; i < shm_size; i++) {
if (i != shm_rank) {
void **ipc_base_uncast = reinterpret_cast<void **>(&ipc_base[i]);
CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i],
@@ -109,7 +109,7 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases,
}
__host__ void IpcOnImpl::ipcHostStop() {
for (size_t i = 0; i < shm_size; i++) {
for (int i = 0; i < shm_size; i++) {
if (i != shm_rank) {
CHECK_HIP(hipIpcCloseMemHandle(ipc_bases[i]));
}
src/ipc_policy.hpp
+1 -1
View File
@@ -44,7 +44,7 @@ class IpcOnImpl {
public:
int shm_rank{0};
uint32_t shm_size{0};
int shm_size{0};
char **ipc_bases{nullptr};
src/util.hpp
+1 -1
View File
@@ -171,7 +171,7 @@ __device__ __forceinline__ void memcpy(void* dst, void* src, size_t size) {
uint8_t* dst_bytes{static_cast<uint8_t*>(dst)};
uint8_t* src_bytes{static_cast<uint8_t*>(src)};
for (int i = 8; i > 1; i >>= 1) {
for (size_t i = 8; i > 1; i >>= 1) {
while (size >= i) {
store_asm(src_bytes, dst_bytes, i);
src_bytes += i;