AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Multi-Node rocshmem_finalize() bug (#138)

Просмотр исходного кода
Этот коммит содержится в:
Yiltan
2025-06-04 10:02:03 -04:00
коммит произвёл GitHub
родитель ca5fdd4718
Коммит 3f01d89207
18 изменённых файлов: 175 добавлений и 237 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
examples/rocshmem_allreduce_test.cc
+5 -18
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
@@ -123,16 +113,13 @@ int main (int argc, char **argv)
nelem = atoi(argv[1]);
}
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int ndevices, my_device = 0;
CHECK_HIP(hipGetDeviceCount(&ndevices));
my_device = my_pe % ndevices;
CHECK_HIP(hipSetDevice(my_device));
CHECK_HIP(hipSetDevice(get_launcher_local_rank()));
rocshmem_init();
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int *source = (int *)rocshmem_malloc(nelem * sizeof(int));
int *dest = (int *)rocshmem_malloc(nelem * sizeof(int));
if (NULL == source || NULL == dest) {
examples/rocshmem_alltoall_test.cc
+5 -18
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
@@ -128,16 +118,13 @@ int main (int argc, char **argv)
nelem = atoi(argv[1]);
}
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int ndevices, my_device = 0;
CHECK_HIP(hipGetDeviceCount(&ndevices));
my_device = my_pe % ndevices;
CHECK_HIP(hipSetDevice(my_device));
CHECK_HIP(hipSetDevice(get_launcher_local_rank()));
rocshmem_init();
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int *source = (int *)rocshmem_malloc(nelem * npes * sizeof(int));
int *dest = (int *)rocshmem_malloc(nelem * npes * sizeof(int));
if (NULL == source || NULL == dest) {
examples/rocshmem_broadcast_test.cc
+5 -18
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
@@ -121,16 +111,13 @@ int main(int argc, char **argv)
nelem = atoi(argv[1]);
}
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int ndevices, my_device = 0;
CHECK_HIP(hipGetDeviceCount(&ndevices));
my_device = my_pe % ndevices;
CHECK_HIP(hipSetDevice(my_device));
CHECK_HIP(hipSetDevice(get_launcher_local_rank()));
rocshmem_init();
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int *source = (int *)rocshmem_malloc(nelem * sizeof(int));
int *dest = (int *)rocshmem_malloc(nelem * sizeof(int));
if (NULL == source || NULL == dest) {
examples/rocshmem_getmem_test.cc
+9 -19
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
@@ -76,8 +66,8 @@ __global__ void simple_getmem_test(int *src, int *dst, size_t nelem)
int threadId = blockIdx.x * blockDim.x + threadIdx.x;
if (threadId == 0) {
int rank = rocshmem_my_pe();
int peer = rank ? 0 : 1;
int my_pe = rocshmem_my_pe();
int peer = my_pe ? 0 : 1;
rocshmem_getmem(dst, src, nelem * sizeof(int), peer);
rocshmem_quiet();
}
@@ -90,19 +80,19 @@ __global__ void simple_getmem_test(int *src, int *dst, size_t nelem)
int main (int argc, char **argv)
{
int rank = rocshmem_my_pe();
int ndevices, my_device = 0;
CHECK_HIP(hipGetDeviceCount(&ndevices));
my_device = rank % ndevices;
CHECK_HIP(hipSetDevice(my_device));
int nelem = MAX_ELEM;
if (argc > 1) {
nelem = atoi(argv[1]);
}
CHECK_HIP(hipSetDevice(get_launcher_local_rank()));
rocshmem_init();
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int *src = (int *)rocshmem_malloc(nelem * sizeof(int));
int *dst = (int *)rocshmem_malloc(nelem * sizeof(int));
if (NULL == src || NULL == dst) {
@@ -128,7 +118,7 @@ int main (int argc, char **argv)
if (dst[i] != 0) {
pass = false;
#if VERBOSE
printf("[%d] Error in element %d expected 0 got %d\n", rank, i, dst[i]);
printf("[%d] Error in element %d expected 0 got %d\n", my_pe, i, dst[i]);
#endif
}
}
examples/rocshmem_init_attr_test.cc
+1 -11
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
examples/rocshmem_put_signal_test.cc
+11 -21
Просмотреть файл
@@ -54,19 +54,9 @@
*/
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#include <rocshmem/rocshmem.hpp>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
#include "util.h"
using namespace rocshmem;
@@ -95,20 +85,20 @@ __global__ void simple_put_signal_test(uint64_t *data, uint64_t *message, size_t
int main (int argc, char **argv)
{
int rank = rocshmem_my_pe();
int ndevices, my_device = 0;
CHECK_HIP(hipGetDeviceCount(&ndevices));
my_device = rank % ndevices;
CHECK_HIP(hipSetDevice(my_device));
int nelem = MAX_ELEM;
if (argc > 1) {
nelem = atoi(argv[1]);
}
CHECK_HIP(hipSetDevice(get_launcher_local_rank()));
rocshmem_init();
int my_pe = rocshmem_my_pe();
int npes = rocshmem_n_pes();
int dst_pe = (rank + 1) % npes;
int dst_pe = (my_pe + 1) % npes;
uint64_t *message = (uint64_t*)rocshmem_malloc(nelem * sizeof(uint64_t));
uint64_t *data = (uint64_t*)rocshmem_malloc(nelem * sizeof(uint64_t));
uint64_t *sig_addr = (uint64_t*)rocshmem_malloc(sizeof(uint64_t));
@@ -123,14 +113,14 @@ int main (int argc, char **argv)
}
for (int i=0; i<nelem; i++) {
message[i] = rank;
message[i] = my_pe;
}
CHECK_HIP(hipMemset(data, 0, (nelem * sizeof(uint64_t))));
CHECK_HIP(hipDeviceSynchronize());
int threadsPerBlock=256;
simple_put_signal_test<<<dim3(1), dim3(threadsPerBlock), 0, 0>>>(data, message, nelem, sig_addr, rank, dst_pe);
simple_put_signal_test<<<dim3(1), dim3(threadsPerBlock), 0, 0>>>(data, message, nelem, sig_addr, my_pe, dst_pe);
rocshmem_barrier_all();
CHECK_HIP(hipDeviceSynchronize());
@@ -139,11 +129,11 @@ int main (int argc, char **argv)
if (data[i] != 0) {
pass = false;
#if VERBOSE
printf("[%d] Error in element %d expected 0 got %d\n", rank, i, dst[i]);
printf("[%d] Error in element %d expected 0 got %d\n", my_pe, i, dst[i]);
#endif
}
}
printf("[%d] Test %s \t %s\n", rank, argv[0], pass ? "[PASS]" : "[FAIL]");
printf("[%d] Test %s \t %s\n", my_pe, argv[0], pass ? "[PASS]" : "[FAIL]");
rocshmem_free(data);
rocshmem_free(message);
examples/util.h
+52
Просмотреть файл
@@ -0,0 +1,52 @@
/******************************************************************************
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*****************************************************************************/
#ifndef __ROCSHMEM_EXAMPLES_UTIL_H__
#define __ROCSHMEM_EXAMPLES_UTIL_H__
#include <iostream>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#define CHECK_HIP(condition) { \
hipError_t error = condition; \
if(error != hipSuccess){ \
fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \
MPI_Abort(MPI_COMM_WORLD, error); \
} \
}
static int get_launcher_local_rank() {
char *local_rank_str = nullptr;
local_rank_str = getenv("OMPI_COMM_WORLD_LOCAL_RANK");
if (nullptr != local_rank_str) {
return atoi(local_rank_str);
}
return -1;
}
#endif /* __ROCSHMEM_EXAMPLES_UTIL_H__ */