AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge branch 'ROCm:develop' into functional_tests/puts_gets

فهرست منبع 
[ROCm/rocshmem commit: e981f61693]
This commit is contained in:
Avinash Kethineedi
2024-10-14 10:27:54 -05:00
کامیت شده توسط GitHub
والد 9532e084fc 5b47cf482d
کامیت 1a5536bfaa
47فایلهای تغییر یافته به همراه2599 افزوده شده و 382 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/rocshmem/.gitignore
+1
مشاهده پرونده
@@ -0,0 +1 @@
build/
projects/rocshmem/CMakeLists.txt
+1 -1
مشاهده پرونده
@@ -61,7 +61,6 @@ option(USE_IPC "Enable IPC support (using HIP)" OFF)
option(USE_THREADS "Enable workgroup threads to share network queues" OFF)
option(USE_WF_COAL "Enable wavefront message coalescing" OFF)
option(USE_COHERENT_HEAP "Enable support for coherent systems" OFF)
option(USE_CACHED_HEAP "Enable support for cached systems" OFF)
option(USE_MANAGED_HEAP "Enable managed memory" OFF)
option(USE_HOST_HEAP "Enable host memory using malloc/free" OFF)
option(USE_HIP_HOST_HEAP "Enable host memory using hip api" OFF)
@@ -69,6 +68,7 @@ option(USE_FUNC_CALL "Force compiler to use function calls on library API" OFF)
option(USE_SHARED_CTX "Request support for shared ctx between WG" OFF)
option(USE_SINGLE_NODE "Enable single node support only." OFF)
option(USE_HOST_SIDE_HDP_FLUSH "Use a polling thread to flush the HDP cache on the host." OFF)
option(USE_COOPERATIVE_GROUPS "Use cooperative groups for internal syncronization" OFF)
option(BUILD_FUNCTIONAL_TESTS "Build the functional tests" ON)
option(BUILD_SOS_TESTS "Build the host-facing tests" OFF)
option(BUILD_UNIT_TESTS "Build the unit tests" ON)
projects/rocshmem/README.md
+1 -1
مشاهده پرونده
@@ -45,7 +45,7 @@ ROC_SHMEM base requirements:
* May work with other versions, but not tested
* AMD GFX9 GPUs (e.g.: MI25, Vega 56, Vega 64, MI50, MI60, MI100, Radeon VII)
* AMD MI200 GPUs: To enable the support on MI200, please configure the library
with USE_CACHED_HEAP
with USE_COHERENT_HEAP
* ROCm-aware MPI as described in
[Building the Dependencies](#building-the-dependencies)
* InfiniBand adaptor compatable with ROCm RDMA technology
projects/rocshmem/cmake/config.h.in
+2 -2
مشاهده پرونده
@@ -8,10 +8,10 @@
#cmakedefine USE_SHARED_CTX
#cmakedefine USE_WF_COAL
#cmakedefine USE_COHERENT_HEAP
#cmakedefine USE_CACHED_HEAP
#cmakedefine USE_MANAGED_HEAP
#cmakedefine USE_HOST_HEAP
#cmakedefine USE_HIP_HOST_HEAP
#cmakedefine USE_FUNC_CALL
#cmakedefine USE_SINGLE_NODE
#cmakedefine USE_HOST_SIDE_HDP_FLUSH
#cmakedefine USE_HOST_SIDE_HDP_FLUSH
#cmakedefine USE_COOPERATIVE_GROUPS
projects/rocshmem/scripts/build_configs/ipc_single_cg
Executable
+31
مشاهده پرونده
@@ -0,0 +1,31 @@
#!/bin/bash
# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
if [ -z $1 ]
then
install_path=~/rocshmem
else
install_path=$1
fi
src_path=$(dirname "$(realpath $0)")/../../
cmake \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=$install_path \
-DCMAKE_VERBOSE_MAKEFILE=OFF \
-DDEBUG=OFF \
-DPROFILE=OFF \
-DUSE_GPU_IB=OFF \
-DUSE_RO=OFF \
-DUSE_DC=OFF \
-DUSE_IPC=ON \
-DUSE_COHERENT_HEAP=ON \
-DUSE_THREADS=OFF \
-DUSE_WF_COAL=OFF \
-DUSE_SINGLE_NODE=ON \
-DUSE_HOST_SIDE_HDP_FLUSH=OFF \
-DUSE_COOPERATIVE_GROUPS=ON \
$src_path
cmake --build . --parallel 8
cmake --install .
projects/rocshmem/scripts/build_configs/rc_single
-1
مشاهده پرونده
@@ -20,7 +20,6 @@ cmake \
-DUSE_DC=OFF \
-DUSE_IPC=OFF \
-DUSE_COHERENT_HEAP=OFF \
-DUSE_CACHED_HEAP=OFF \
-DUSE_THREADS=OFF \
-DUSE_WF_COAL=OFF \
$src_path
projects/rocshmem/scripts/build_configs/rc_single_single_node
-2
مشاهده پرونده
@@ -20,12 +20,10 @@ cmake \
-DUSE_DC=OFF \
-DUSE_IPC=ON \
-DUSE_COHERENT_HEAP=OFF \
-DUSE_CACHED_HEAP=OFF \
-DUSE_THREADS=OFF \
-DUSE_WF_COAL=OFF \
-DUSE_SINGLE_NODE=ON \
-DUSE_HOST_SIDE_HDP_FLUSH=ON\
-DROCM_PATH="/opt/rocm-5.4.2/"\
$src_path
cmake --build . --parallel 8
cmake --install .
projects/rocshmem/scripts/build_configs/rc_single_single_node_debug
-1
مشاهده پرونده
@@ -20,7 +20,6 @@ cmake \
-DUSE_DC=OFF \
-DUSE_IPC=ON \
-DUSE_COHERENT_HEAP=OFF \
-DUSE_CACHED_HEAP=OFF \
-DUSE_THREADS=OFF \
-DUSE_WF_COAL=OFF \
-DUSE_SINGLE_NODE=ON \
projects/rocshmem/scripts/build_configs/ro_ipc
Executable
+29
مشاهده پرونده
@@ -0,0 +1,29 @@
#!/bin/bash
# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
if [ -z $1 ]
then
install_path=~/rocshmem
else
install_path=$1
fi
src_path=$(dirname "$(realpath $0)")/../../
cmake \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=$install_path \
-DCMAKE_VERBOSE_MAKEFILE=OFF \
-DDEBUG=OFF \
-DPROFILE=OFF \
-DUSE_GPU_IB=OFF \
-DUSE_DC=OFF \
-DUSE_IPC=ON \
-DUSE_COHERENT_HEAP=ON \
-DUSE_THREADS=OFF \
-DUSE_WF_COAL=OFF \
-DUSE_SINGLE_NODE=ON \
-DUSE_HOST_SIDE_HDP_FLUSH=OFF\
$src_path
cmake --build . --parallel 8
cmake --install .
projects/rocshmem/scripts/build_configs/ro_net
+1 -1
مشاهده پرونده
@@ -22,7 +22,7 @@ cmake \
-DUSE_IPC=OFF \
-DUSE_THREADS=ON \
-DUSE_WF_COAL=OFF \
-DUSE_COHERENT_HEAP=ON \
-DUSE_COHERENT_HEAP=OFF \
$src_path
cmake --build . --parallel 8
cmake --install .
projects/rocshmem/scripts/build_configs/ro_net_debug
+1 -1
مشاهده پرونده
@@ -21,7 +21,7 @@ cmake \
-DUSE_IPC=OFF \
-DUSE_THREADS=ON \
-DUSE_WF_COAL=OFF \
-DUSE_COHERENT_HEAP=ON \
-DUSE_COHERENT_HEAP=OFF \
$src_path
cmake --build . --parallel 8
cmake --install .
projects/rocshmem/scripts/functional_tests/driver.sh
+6 -6
مشاهده پرونده
@@ -121,9 +121,9 @@ case $2 in
echo "amoinc_n2_w1_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -a 11 > $3/amoinc_n2_w1_z1.log
check amoinc_n2_w1_z1
echo "pingpong_n2_w1"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log
check pingpong_n2_w1
# echo "pingpong_n2_w1"
# ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log
# check pingpong_n2_w1
echo "amoset_n2_w1_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -a 44 > $3/amoset_n2_w1_z1.log
check amoset_n2_w1_z1
@@ -163,9 +163,9 @@ case $2 in
echo "amoinc_n2_w8_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 11 > $3/amoinc_n2_w8_z1.log
check amoinc_n2_w8_z1
echo "pingpong_n2_w1"
ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log
check pingpong_n2_w1
# echo "pingpong_n2_w1"
# ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log
# check pingpong_n2_w1
echo "amoset_n2_w8_z1"
ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 44 > $3/amoset_n2_w8_z1.log
check amoset_n2_w8_z1
projects/rocshmem/scripts/unit_tests/driver.sh
Executable
+62
مشاهده پرونده
@@ -0,0 +1,62 @@
#!/bin/bash
# Function to display help information
function display_help {
echo "Usage:"
echo " $0 binary_name all # Runs all standard tests"
echo " $0 binary_name custom <ranks> <filter> # Runs custom test configuration"
echo
echo "Arguments:"
echo " binary_name: Name of the binary to run."
echo " all: Executes predefined test configurations."
echo " custom: Executes a test with custom MPI ranks and GTest filter."
echo " ranks: Number of MPI ranks (required for custom mode)."
echo " filter: GTest filter string (required for custom mode)."
echo
}
# Validate number of arguments for each mode
if [[ "$#" -lt 2 ]] ||
{ [[ "$2" == "all" ]] && [[ "$#" -ne 2 ]]; } ||
{ [[ "$2" == "custom" ]] && [[ "$#" -ne 4 ]]; }; then
display_help
exit 1
fi
binary_name=$1
mode=$2
timestamp=$(date "+%Y-%m-%d-%H:%M:%S")
log_file="unit_tests_${timestamp}.log"
# Function to execute mpirun command
function run_mpirun {
local np=$1
local gtest_filter=$2
echo "mpirun -np $np $binary_name --gtest_filter='$gtest_filter'" | tee -a "$log_file"
mpirun -np "$np" "$binary_name" --gtest_filter="$gtest_filter" >> "$log_file" 2>&1
}
# Processing modes
case $mode in
all)
run_mpirun 4 "-IPCImplSimpleCoarseTestFixture.*:IPCImplSimpleFineTestFixture.*"
run_mpirun 2 "IPCImplSimpleCoarseTestFixture.*:IPCImplSimpleFineTestFixture.*"
;;
custom)
# Check if ranks is a positive integer
if [[ "$3" -le 1 ]]; then
echo "Error: 'ranks' must be a positive integer."
display_help
exit 1
fi
run_mpirun $3 $4
;;
*)
echo "Error: Invalid mode '$mode'." | tee -a "$log_file"
display_help
exit 1
;;
esac
echo "Tests Completed"
echo "log file: '$log_file'"
projects/rocshmem/src/assembly.hpp
+271
مشاهده پرونده
@@ -0,0 +1,271 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* 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 LIBRARY_SRC_ASSEMBLY_HPP_
#define LIBRARY_SRC_ASSEMBLY_HPP_
#include <hip/hip_runtime.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
namespace rocshmem {
#define DO_PRAGMA(x) _Pragma(#x)
#define NOWARN(warnoption, ...) \
DO_PRAGMA(GCC diagnostic push) \
DO_PRAGMA(GCC diagnostic ignored #warnoption) \
__VA_ARGS__ \
DO_PRAGMA(GCC diagnostic pop)
#define SFENCE() asm volatile("sfence" ::: "memory")
__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
int ret;
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"global_load_ubyte %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"global_load_ubyte %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
return ret;
}
__device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_value,
volatile char *read_value) {
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"global_load_sbyte %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"global_load_sbyte %0 %1 off sc0 sc1\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
#endif
}
__device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value,
volatile uint64_t *read_value) {
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"global_load_dwordx2 %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"global_load_dwordx2 %0 %1 off sc0 sc1\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
#endif
}
/* Ignore the warning about deprecated volatile.
* The only usage of volatile is to force the compiler to generate
* the assembly instruction. If volatile is omitted, the compiler
* will NOT generate the non-temporal load or the waitcnt.
*/
// clang-format off
NOWARN(-Wdeprecated-volatile,
template <typename T> __device__ __forceinline__ T uncached_load(T* src) {
T ret;
switch (sizeof(T)) {
case 4:
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"global_load_dword %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"global_load_dword %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
break;
case 8:
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"global_load_dwordx2 %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"global_load_dwordx2 %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
#endif
break;
default:
break;
}
return ret;
}
)
// clang-format on
__device__ __forceinline__ void __roc_inv() {
#if defined USE_COHERENT_HEAP
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
// asm volatile("buffer_wbinvl1;");
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
// asm volatile("buffer_inv sc0 sc1;");
#endif
#endif
}
__device__ __forceinline__ void __roc_flush() {
#if defined USE_COHERENT_HEAP
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
// asm volatile("s_dcache_wb;");
// asm volatile("buffer_wbl2;");
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
// asm volatile("s_dcache_wb;");
// asm volatile("buffer_wbl2;");
#endif
#endif
}
__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
int size) {
switch (size) {
case 2: {
int16_t val16{*(reinterpret_cast<int16_t*>(val))};
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16));
#endif
break;
}
case 4: {
int32_t val32{*(reinterpret_cast<int32_t*>(val))};
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32));
#endif
break;
}
case 8: {
int64_t val64{*(reinterpret_cast<int64_t*>(val))};
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64));
#endif
break;
}
default:
break;
}
}
__device__ __forceinline__ uint64_t __read_clock() {
uint64_t clock{};
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
asm volatile(
"s_memrealtime %0\n"
"s_waitcnt lgkmcnt(0)\n"
: "=s"(clock));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
asm volatile(
"s_memrealtime %0\n"
"s_waitcnt lgkmcnt(0)\n"
: "=s"(clock));
#endif
return clock;
}
} // namespace rocshmem
#endif // LIBRARY_SRC_ASSEMBLY_HPP_
projects/rocshmem/src/atomic.hpp
+141
مشاهده پرونده
@@ -0,0 +1,141 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* 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 LIBRARY_SRC_ATOMIC_HPP
#define LIBRARY_SRC_ATOMIC_HPP
#include <hip/hip_runtime.h>
namespace rocshmem {
namespace detail {
namespace atomic {
typedef enum rocshmem_memory_scope {
memory_scope_thread = __HIP_MEMORY_SCOPE_SINGLETHREAD,
memory_scope_wavefront = __HIP_MEMORY_SCOPE_WAVEFRONT,
memory_scope_workgroup = __HIP_MEMORY_SCOPE_WORKGROUP,
memory_scope_agent = __HIP_MEMORY_SCOPE_AGENT,
memory_scope_system = __HIP_MEMORY_SCOPE_SYSTEM,
} rocshmem_memory_scope;
typedef enum rocshmem_memory_order {
memory_order_relaxed = __ATOMIC_RELAXED,
memory_order_consume = __ATOMIC_CONSUME,
memory_order_acquire = __ATOMIC_ACQUIRE,
memory_order_release = __ATOMIC_RELEASE,
memory_order_acq_rel = __ATOMIC_ACQ_REL,
memory_order_seq_cst = __ATOMIC_SEQ_CST
} rocshmem_memory_order;
struct rocshmem_memory_orders {
rocshmem_memory_order load {memory_order_acquire};
rocshmem_memory_order store {memory_order_release};
rocshmem_memory_order atomic {memory_order_acq_rel};
rocshmem_memory_order weak_cas_success {memory_order_acq_rel};
rocshmem_memory_order weak_cas_failure {memory_order_acq_rel};
rocshmem_memory_order strong_cas_success {memory_order_acq_rel};
rocshmem_memory_order strong_cas_failure {memory_order_acq_rel};
};
template <typename T, rocshmem_memory_scope s>
__host__ __device__
T load(const T* address, rocshmem_memory_orders o) {
return __hip_atomic_load(address, o.load, s);
}
template <typename T, rocshmem_memory_scope s>
__host__ __device__
void store(T* address, const T value, rocshmem_memory_orders o) {
return __hip_atomic_store(address, value, o.store, s);
}
template <typename T, rocshmem_memory_scope s>
__host__ __device__
bool compare_exchange_weak(T& expected, T desired, rocshmem_memory_orders o) {
return __hip_atomic_compare_exchange_weak(expected, desired, o.weak_cas_success, o.weak_cas_failure, s);
}
template <typename T, rocshmem_memory_scope s>
__host__ __device__
bool compare_exchange_strong(T& expected, T desired, rocshmem_memory_orders o) {
return __hip_atomic_compare_exchange_strong(expected, desired, o.strong_cas_success, o.strong_cas_failure, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_add(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_add(obj, arg, o.atomic, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_sub(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_sub(obj, arg, o.atomic, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_and(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_and(obj, arg, o.atomic, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_or(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_or(obj, arg, o, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_xor(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_xor(obj, arg, o.atomic, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_max(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_max(obj, arg, o.atomic, s);
}
template <class T, class U, rocshmem_memory_scope s>
__host__ __device__
T fetch_min(T* obj, U arg, rocshmem_memory_orders o) {
return __hip_atomic_fetch_min(obj, arg, o.atomic, s);
}
template <rocshmem_memory_scope s>
__device__
void threadfence() {
if constexpr (s == memory_scope_system) {
__threadfence_system();
} else if constexpr (s == memory_scope_agent) {
__threadfence();
} else if constexpr (s == memory_scope_workgroup) {
__threadfence_block();
}
}
} // namespace atomic
} // namespace detail
} // namespace rocshmem
#endif // LIBRARY_SRC_ATOMIC_HPP_
projects/rocshmem/src/backend_type.hpp
+4 -8
مشاهده پرونده
@@ -70,13 +70,11 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND };
*/
#ifdef USE_GPU_IB
#define DISPATCH_RET(Func) \
auto ret_val{0}; \
ret_val = static_cast<GPUIBContext *>(this)->Func; \
auto ret_val = static_cast<GPUIBContext *>(this)->Func; \
return ret_val;
#elif defined(USE_RO)
#define DISPATCH_RET(Func) \
auto ret_val{0}; \
ret_val = static_cast<ROContext *>(this)->Func; \
auto ret_val = static_cast<ROContext *>(this)->Func; \
return ret_val;
#else
#define DISPATCH_RET(Func) \
@@ -128,13 +126,11 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND };
#ifdef USE_GPU_IB
#define HOST_DISPATCH_RET(Func) \
auto ret_val{0}; \
ret_val = static_cast<GPUIBHostContext *>(this)->Func; \
auto ret_val = static_cast<GPUIBHostContext *>(this)->Func; \
return ret_val;
#elif defined(USE_RO)
#define HOST_DISPATCH_RET(Func) \
auto ret_val{0}; \
ret_val = static_cast<ROHostContext *>(this)->Func; \
auto ret_val = static_cast<ROHostContext *>(this)->Func; \
return ret_val;
#else
#define HOST_DISPATCH_RET(Func) \
projects/rocshmem/src/gpu_ib/backend_ib.cpp
+2 -3
مشاهده پرونده
@@ -133,8 +133,6 @@ __device__ bool GPUIBBackend::create_ctx(int64_t options,
return false;
}
ctx_ = pop_result.value;
ctx_->dev_mtx_.shareable_ = static_cast<bool>(options >> 3);
ctx->ctx_opaque = ctx_;
return true;
}
@@ -296,7 +294,8 @@ void GPUIBBackend::setup_default_ctx() {
CHECK_HIP(hipGetSymbolAddress(reinterpret_cast<void **>(&symbol_address),
HIP_SYMBOL(ROC_SHMEM_CTX_DEFAULT)));
roc_shmem_ctx_t ctx_default_host{default_ctx_, nullptr};
TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world;
roc_shmem_ctx_t ctx_default_host{default_ctx_, tinfo};
hipStream_t stream;
CHECK_HIP(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
projects/rocshmem/src/gpu_ib/connection.cpp
+1 -1
مشاهده پرونده
@@ -256,7 +256,7 @@ void* Connection::buf_alloc([[maybe_unused]] struct ibv_pd* pd,
if (use_gpu_mem) {
void* dev_ptr;
if (coherent_cq == 1) {
#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
#if defined USE_COHERENT_HEAP
CHECK_HIP(hipMalloc(reinterpret_cast<void**>(&dev_ptr), size));
#else
#ifdef HIP_SUPPORTS_MALLOC_UNCACHED
projects/rocshmem/src/gpu_ib/segment_builder.cpp
+30 -22
مشاهده پرونده
@@ -22,6 +22,7 @@
#include "segment_builder.hpp"
#include "../util.hpp"
#include "endian.hpp"
namespace rocshmem {
@@ -35,9 +36,9 @@ __device__ SegmentBuilder::SegmentBuilder(uint64_t wqe_idx, void *base) {
__device__ void SegmentBuilder::update_cntrl_seg(
uint8_t opcode, uint16_t wqe_idx, uint32_t ctrl_qp_sq, uint64_t ctrl_sig,
ConnectionImpl *connection_policy, bool zero_byte_rd) {
mlx5_wqe_ctrl_seg *ctrl_seg = &seg_ptr->ctrl_seg;
mlx5_wqe_ctrl_seg ctrl_seg;
ctrl_seg->opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8);
ctrl_seg.opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8);
uint32_t DS = 2;
if (zero_byte_rd == false) {
@@ -48,39 +49,43 @@ __device__ void SegmentBuilder::update_cntrl_seg(
DS += connection_policy->wqeCntrlOffset();
ctrl_seg->qpn_ds = (DS << 24) | ctrl_qp_sq;
ctrl_seg.qpn_ds = (DS << 24) | ctrl_qp_sq;
ctrl_seg->signature = ctrl_sig;
ctrl_seg.signature = ctrl_sig;
ctrl_seg->fm_ce_se = ctrl_sig >> 24;
ctrl_seg.fm_ce_se = ctrl_sig >> 24;
ctrl_seg->imm = ctrl_sig >> 32;
ctrl_seg.imm = ctrl_sig >> 32;
memcpy(&seg_ptr->ctrl_seg, &ctrl_seg, sizeof(mlx5_wqe_ctrl_seg));
seg_ptr++;
}
__device__ void SegmentBuilder::update_atomic_data_seg(uint64_t atomic_data,
uint64_t atomic_cmp) {
mlx5_wqe_atomic_seg *atomic_seg = &seg_ptr->atomic_seg;
mlx5_wqe_atomic_seg atomic_seg;
swap_endian_store(reinterpret_cast<uint64_t *>(&atomic_seg->swap_add),
swap_endian_store(reinterpret_cast<uint64_t *>(&atomic_seg.swap_add),
atomic_data);
swap_endian_store(reinterpret_cast<uint64_t *>(&atomic_seg->compare),
swap_endian_store(reinterpret_cast<uint64_t *>(&atomic_seg.compare),
atomic_cmp);
memcpy(&seg_ptr->atomic_seg, &atomic_seg, sizeof(mlx5_wqe_atomic_seg));
seg_ptr++;
}
__device__ void SegmentBuilder::update_rdma_seg(uintptr_t *raddr,
uint32_t rkey) {
mlx5_wqe_raddr_seg *raddr_seg = &seg_ptr->raddr_seg;
mlx5_wqe_raddr_seg raddr_seg;
raddr_seg->rkey = rkey;
raddr_seg.rkey = rkey;
swap_endian_store(reinterpret_cast<uint64_t *>(&raddr_seg->raddr),
swap_endian_store(reinterpret_cast<uint64_t *>(&raddr_seg.raddr),
reinterpret_cast<uint64_t>(raddr));
memcpy(&seg_ptr->raddr_seg, &raddr_seg, sizeof(mlx5_wqe_raddr_seg));
seg_ptr++;
}
@@ -90,33 +95,36 @@ __device__ void SegmentBuilder::update_data_seg(uintptr_t *laddr, int32_t size,
return;
}
mlx5_wqe_data_seg *data_seg = &seg_ptr->data_seg;
mlx5_wqe_data_seg data_seg;
data_seg.lkey = lkey;
data_seg->lkey = lkey;
swap_endian_store(&data_seg->byte_count, size & 0x7FFFFFFFU);
swap_endian_store(reinterpret_cast<uint64_t *>(&data_seg->addr),
swap_endian_store(&data_seg.byte_count, size & 0x7FFFFFFFU);
swap_endian_store(reinterpret_cast<uint64_t *>(&data_seg.addr),
reinterpret_cast<uint64_t>(laddr));
memcpy(&seg_ptr->data_seg, &data_seg, sizeof(mlx5_wqe_data_seg));
seg_ptr++;
}
__device__ void SegmentBuilder::update_inl_data_seg(uintptr_t *laddr,
int32_t size) {
mlx5_wqe_inl_data_seg *inl_data_seg = &seg_ptr->inl_data_seg;
mlx5_wqe_inl_data_seg inl_data_seg;
swap_endian_store(&inl_data_seg->byte_count, (size & 0x3FF) | 0x80000000);
swap_endian_store(&inl_data_seg.byte_count, (size & 0x3FF) | 0x80000000);
// Assume fence HDP flush
// TODO(khamidou): Rework fence interface to avoid this
size_t field_size{sizeof(mlx5_wqe_inl_data_seg)};
if (!laddr) {
uint8_t flush_val = 1;
memcpy(inl_data_seg + 1, &flush_val, sizeof(flush_val));
memcpy(&inl_data_seg + 1, &flush_val, sizeof(flush_val));
field_size += sizeof(flush_val);
} else {
memcpy(inl_data_seg + 1, laddr, size);
memcpy(&inl_data_seg + 1, laddr, size);
field_size += size;
}
memcpy(&seg_ptr->inl_data_seg, &inl_data_seg, field_size);
seg_ptr++;
}
projects/rocshmem/src/ipc/CMakeLists.txt
+1
مشاهده پرونده
@@ -30,4 +30,5 @@ target_sources(
context_ipc_host.cpp
backend_ipc.cpp
ipc_team.cpp
context_ipc_device_coll.cpp
)
projects/rocshmem/src/ipc/backend_ipc.cpp
+22 -9
مشاهده پرونده
@@ -56,18 +56,18 @@ int get_ls_non_zero_bit(char *bitmask, int mask_length) {
IPCBackend::IPCBackend(MPI_Comm comm)
: Backend() {
type = BackendType::IPC_BACKEND;
if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) {
std::stringstream sstream(maximum_num_contexts_str);
sstream >> maximum_num_contexts_;
}
init_mpi_once(comm);
initIPC();
auto *bp{ipc_backend_proxy.get()};
bp->heap_ptr = &heap;
/* Initialize the host interface */
@@ -82,16 +82,19 @@ IPCBackend::IPCBackend(MPI_Comm comm)
allocate_atomic_region(&bp->atomic_ret, MAX_NUM_BLOCKS);
default_context_proxy_ = IPCDefaultContextProxyT(this);
setup_team_world();
TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world;
roc_shmem_collective_init();
setup_fence_buffer();
teams_init();
default_context_proxy_ = IPCDefaultContextProxyT(this, tinfo);
setup_ctxs();
}
IPCBackend::~IPCBackend() {
@@ -105,7 +108,7 @@ IPCBackend::~IPCBackend() {
* Free the atomic_ret array.
*/
CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr));
// TODO(Avinash) Free g_ret
// delete host_interface;
@@ -141,6 +144,8 @@ __device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
ctx_ = pop_result.value;
ctx->ctx_opaque = ctx_;
ctx_->tinfo = reinterpret_cast<TeamInfo *>(ctx->team_opaque);
return true;
}
@@ -287,6 +292,14 @@ void IPCBackend::teams_destroy() {
free(reduced_bitmask_);
}
void IPCBackend::setup_fence_buffer() {
/*
* Allocate heap space for fence
*/
fence_pool = reinterpret_cast<int *>(roc_shmem_malloc(
sizeof(int) * num_pes));
}
void IPCBackend::roc_shmem_collective_init() {
/*
* Allocate heap space for barrier_sync
@@ -389,4 +402,4 @@ void IPCBackend::teams_init() {
NET_CHECK(MPI_Barrier(thread_comm));
}
} // namespace rocshmem
} // namespace rocshmem
projects/rocshmem/src/ipc/backend_ipc.hpp
+20 -2
مشاهده پرونده
@@ -32,6 +32,10 @@
#include "ipc_context_proxy.hpp"
#include "../ipc_policy.hpp"
#ifndef USE_COOPERATIVE_GROUPS
#include "../memory/notifier.hpp"
#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
namespace rocshmem {
class IPCBackend : public Backend {
@@ -104,7 +108,7 @@ class IPCBackend : public Backend {
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) override;
/**
* @copydoc Backend::team_destroy(roc_shmem_team_t)
*/
@@ -156,6 +160,15 @@ class IPCBackend : public Backend {
*/
void *pAta_pool{nullptr};
/**
* @brief Handle for raw memory for fence/quiet
*/
int *fence_pool{nullptr};
#ifndef USE_COOPERATIVE_GROUPS
NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent> notifier_{};
#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
protected:
/**
* @copydoc Backend::dump_backend_stats()
@@ -203,6 +216,11 @@ class IPCBackend : public Backend {
*/
void roc_shmem_collective_init();
/**
* @brief Allocate buffer for fence/quiet operation
*/
void setup_fence_buffer();
private:
/**
* @brief Proxy for the default context
@@ -249,4 +267,4 @@ class IPCBackend : public Backend {
} // namespace rocshmem
#endif // LIBRARY_SRC_IPC_BACKEND_HPP_
#endif // LIBRARY_SRC_IPC_BACKEND_HPP_
projects/rocshmem/src/ipc/context_ipc_device.cpp
+13 -12
مشاهده پرونده
@@ -44,8 +44,16 @@ __host__ IPCContext::IPCContext(Backend *b)
auto *bp{backend->ipc_backend_proxy.get()};
barrier_sync = backend->barrier_sync;
g_ret = bp->g_ret;
atomic_base_ptr = bp->atomic_ret->atomic_base_ptr;
fence_pool = backend->fence_pool;
orders_.store = detail::atomic::rocshmem_memory_order::memory_order_seq_cst;
#ifndef USE_COOPERATIVE_GROUPS
notifier_ = backend->notifier_.get();
#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
}
__device__ void IPCContext::threadfence_system() {
@@ -84,12 +92,17 @@ __device__ void IPCContext::getmem_nbi(void *dest, const void *source,
}
__device__ void IPCContext::fence() {
for (int i{0}, j{tinfo->pe_start}; i < tinfo->size; i++, j += tinfo->stride) {
detail::atomic::store<int, detail::atomic::memory_scope_system>(&fence_pool[j], 1, orders_);
}
}
__device__ void IPCContext::fence(int pe) {
detail::atomic::store<int, detail::atomic::memory_scope_system>(&fence_pool[pe], 1, orders_);
}
__device__ void IPCContext::quiet() {
fence();
}
__device__ void *IPCContext::shmem_ptr(const void *dest, int pe) {
@@ -97,18 +110,6 @@ __device__ void *IPCContext::shmem_ptr(const void *dest, int pe) {
return ret;
}
__device__ void IPCContext::barrier_all() {
__syncthreads();
}
__device__ void IPCContext::sync_all() {
__syncthreads();
}
__device__ void IPCContext::sync(roc_shmem_team_t team) {
__syncthreads();
}
__device__ void IPCContext::putmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
uint64_t L_offset =
projects/rocshmem/src/ipc/context_ipc_device.hpp
+61 -40
مشاهده پرونده
@@ -24,6 +24,15 @@
#define LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_
#include "../context.hpp"
#include "../atomic.hpp"
#include "../team.hpp"
#ifdef USE_COOPERATIVE_GROUPS
#include <hip/hip_cooperative_groups.h>
namespace cg = cooperative_groups;
#else
#include "../memory/notifier.hpp"
#endif /* USE_COOPERATIVE_GROUPS */
namespace rocshmem {
@@ -194,50 +203,62 @@ class IPCContext : public Context {
template <typename T>
__device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
// Wait / Test functions
template <typename T>
__device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val);
template <typename T>
__device__ void wait_until_all(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val);
template <typename T>
__device__ size_t wait_until_any(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val);
template <typename T>
__device__ size_t wait_until_some(T* ptr, size_t nelems,
size_t* indices,
const int *status,
roc_shmem_cmps cmp, T val);
template <typename T>
__device__ void wait_until_all_vector(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T* vals);
template <typename T>
__device__ size_t wait_until_any_vector(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T* vals);
template <typename T>
__device__ size_t wait_until_some_vector(T* ptr, size_t nelems,
size_t* indices,
const int *status,
roc_shmem_cmps cmp, T* vals);
template <typename T>
__device__ int test(T* ptr, roc_shmem_cmps cmp, T val);
private:
uint64_t* atomic_base_ptr{nullptr};
//context class has IpcImpl object (ipcImpl_)
IpcImpl *ipcImpl{nullptr};
char* g_ret;
uint64_t* atomic_base_ptr{nullptr};
char* g_ret;
//internal functions used by collective operations
template <typename T>
__device__ void internal_put_broadcast(T *dst, const T *src, int nelems,
int pe_root, int PE_start,
int logPE_stride, int PE_size); // NOLINT(runtime/int)
template <typename T>
__device__ void internal_get_broadcast(T *dst, const T *src, int nelems,
int pe_root); // NOLINT(runtime/int)
template <typename T>
__device__ void fcollect_linear(roc_shmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void alltoall_linear(roc_shmem_team_t team, T *dest,
const T *source, int nelems);
__device__ void internal_sync(int pe, int PE_start, int stride, int PE_size,
int64_t *pSync);
__device__ void internal_direct_barrier(int pe, int PE_start, int stride,
int n_pes, int64_t *pSync);
__device__ void internal_atomic_barrier(int pe, int PE_start, int stride,
int n_pes, int64_t *pSync);
//Temporary scratchpad memory used by internal barrier algorithms.
int64_t *barrier_sync{nullptr};
#ifndef USE_COOPERATIVE_GROUPS
Notifier<detail::atomic::memory_scope_agent> *notifier_{nullptr};
#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */
//Struct defining memory ordering for atomic operations.
detail::atomic::rocshmem_memory_orders orders_{};
//Buffer to perform Atomic store to enforce memory ordering
int *fence_pool{nullptr};
public:
//TODO(Avinash):
//Make tinfo private variable, it requires changes to the context
//creation API in backend
//Team information for the team associated with the context
TeamInfo *tinfo{nullptr};
};
} // namespace rocshmem
projects/rocshmem/src/ipc/context_ipc_device_coll.cpp
+137
مشاهده پرونده
@@ -0,0 +1,137 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* 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.
*****************************************************************************/
#include "roc_shmem/roc_shmem.hpp"
#include "../context_incl.hpp"
#include "context_ipc_tmpl_device.hpp"
#include "../util.hpp"
#include "ipc_team.hpp"
namespace rocshmem {
__device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
int stride, int n_pes,
int64_t *pSync) {
int64_t flag_val = 1;
if (pe == PE_start) {
// Go through all PE offsets (except current offset = 0)
// and wait until they all reach
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
for (size_t i = 1; i < n_pes; i++) {
wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
pSync[i] = ROC_SHMEM_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) {
put_nbi(&pSync[0], &flag_val, 1, j);
}
} else {
// Mark current PE offset as reached
size_t pe_offset = (pe - PE_start) / stride;
put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start);
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
threadfence_system();
}
}
__device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
int stride, int n_pes,
int64_t *pSync) {
int64_t flag_val = 1;
if (pe == PE_start) {
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1));
pSync[0] = ROC_SHMEM_SYNC_VALUE;
threadfence_system();
for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
put_nbi(&pSync[0], &flag_val, 1, j);
}
} else {
amo_add<int64_t>(&pSync[0], flag_val, PE_start);
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
threadfence_system();
}
}
// Uses PE values that are relative to world
__device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
int PE_size, int64_t *pSync) {
#ifdef USE_COOPERATIVE_GROUPS
cg::grid_group grid = cg::this_grid();
grid.sync();
if (0 == grid.thread_rank())
#else
notifier_->sync();
if (0 == get_flat_id())
#endif /* USE_COOPERATIVE_GROUPS */
{
if (PE_size < 64) {
internal_direct_barrier(pe, PE_start, stride, PE_size, pSync);
} else {
internal_atomic_barrier(pe, PE_start, stride, PE_size, pSync);
}
}
__threadfence();
#ifdef USE_COOPERATIVE_GROUPS
grid.sync();
#else
notifier_->sync();
#endif /* USE_COOPERATIVE_GROUPS */
}
__device__ void IPCContext::sync(roc_shmem_team_t team) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
/**
* Ensure that the stride is a multiple of 2.
*/
int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
int pe = team_obj->my_pe_in_world;
int pe_start = team_obj->tinfo_wrt_world->pe_start;
int pe_stride = (1 << log_pe_stride);
int pe_size = team_obj->num_pes;
internal_sync(pe, pe_start, pe_stride, pe_size, barrier_sync);
}
__device__ void IPCContext::sync_all() {
internal_sync(my_pe, 0, 1, num_pes, barrier_sync);
}
__device__ void IPCContext::barrier_all() {
if (is_thread_zero_in_block()) {
quiet();
}
sync_all();
__syncthreads();
}
} // namespace rocshmem
projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp
+104 -54
مشاهده پرونده
@@ -27,6 +27,7 @@
#include "roc_shmem/roc_shmem.hpp"
#include "context_ipc_device.hpp"
#include "../util.hpp"
#include "ipc_team.hpp"
namespace rocshmem {
@@ -165,11 +166,46 @@ __device__ void IPCContext::to_all(T *dest, const T *source, int nreduce,
long *pSync) { // NOLINT(runtime/int)
}
template <typename T>
__device__ void IPCContext::internal_put_broadcast(
T *dst, const T *src, int nelems, int pe_root, int pe_start,
int log_pe_stride, int pe_size) { // NOLINT(runtime/int)
if (my_pe == pe_root) {
int stride = 1 << log_pe_stride;
int finish = pe_start + stride * pe_size;
for (int i = pe_start; i < finish; i += stride) {
if (i != my_pe) {
put_nbi_wg(dst, src, nelems, i);
}
}
}
}
template <typename T>
__device__ void IPCContext::internal_get_broadcast(
T *dst, const T *src, int nelems, int pe_root) { // NOLINT(runtime/int)
if (my_pe != pe_root) {
get_wg(dst, src, nelems, pe_root);
}
}
template <typename T>
__device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst,
const T *src, int nelems, int pe_root) {
//broadcast<T>(dst, src, nelems, pe_root_world, pe_start, log_pe_stride,
// pe_size, p_sync);
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
/**
* Ensure that the stride is a multiple of 2 .
*/
int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
int pe_start = team_obj->tinfo_wrt_world->pe_start;
int pe_size = team_obj->tinfo_wrt_world->size;
long *p_sync = team_obj->bcast_pSync;
// Passed pe_root is relative to team, convert to world root
int pe_root_world = team_obj->get_pe_in_world(pe_root);
broadcast<T>(dst, src, nelems, pe_root_world, pe_start, log_pe_stride,
pe_size, p_sync);
}
template <typename T>
@@ -177,16 +213,82 @@ __device__ void IPCContext::broadcast(T *dst, const T *src, int nelems,
int pe_root, int pe_start,
int log_pe_stride, int pe_size,
long *p_sync) { // NOLINT(runtime/int)
if (num_pes < 4) {
internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride,
pe_size);
} else {
internal_get_broadcast(dst, src, nelems, pe_root);
}
// Synchronize on completion of broadcast
internal_sync(my_pe, pe_start, (1 << log_pe_stride), pe_size, p_sync);
}
template <typename T>
__device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst,
const T *src, int nelems) {
alltoall_linear(team, dst, src, nelems);
}
template <typename T>
__device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst,
const T *src, int nelems) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
/**
* Ensure that the stride is a multiple of 2
*/
int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
int pe_start = team_obj->tinfo_wrt_world->pe_start;
int pe_size = team_obj->num_pes;
int stride = 1 << log_pe_stride;
long *pSync = team_obj->alltoall_pSync;
int my_pe_in_team = team_obj->my_pe;
// Have each PE put their designated data to the other PEs
for (int j = 0; j < pe_size; j++) {
int dest_pe = team_obj->get_pe_in_world(j);
put_nbi_wg(&dst[my_pe_in_team * nelems], &src[j * nelems], nelems, dest_pe);
}
if (is_thread_zero_in_block()) {
quiet();
}
// wait until everyone has obtained their designated data
internal_sync(my_pe, pe_start, stride, pe_size, pSync);
}
template <typename T>
__device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst,
const T *src, int nelems) {
fcollect_linear(team, dst, src, nelems);
}
template <typename T>
__device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst,
const T *src, int nelems) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
/**
* Ensure that the stride is a multiple of 2.
*/
int log_pe_stride = static_cast<int>(team_obj->tinfo_wrt_world->log_stride);
int pe_start = team_obj->tinfo_wrt_world->pe_start;
int pe_size = team_obj->num_pes;
int stride = 1 << log_pe_stride;
long *pSync = team_obj->alltoall_pSync;
int my_pe_in_team = team_obj->my_pe;
// Have each PE put their designated data to the other PEs
for (int j = 0; j < pe_size; j++) {
int dest_pe = team_obj->get_pe_in_world(j);
put_nbi_wg(&dst[my_pe_in_team * nelems], src, nelems, dest_pe);
}
if (is_thread_zero_in_block()) {
quiet();
}
// wait until everyone has obtained their designated data
internal_sync(my_pe, pe_start, stride, pe_size, pSync);
}
// Block/wave functions
@@ -238,58 +340,6 @@ __device__ void IPCContext::get_nbi_wave(T *dest, const T *source,
getmem_nbi_wave(dest, source, nelems * sizeof(T), pe);
}
//Wait/test functions
template <typename T>
__device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val) {
}
template <typename T>
__device__ void wait_until_all(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val) {
}
template <typename T>
__device__ size_t wait_until_any(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val) {
return 0;
}
template <typename T>
__device__ size_t wait_until_some(T* ptr, size_t nelems,
size_t* indices,
const int *status,
roc_shmem_cmps cmp, T val){
return 0;
}
template <typename T>
__device__ void wait_until_all_vector(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T* vals) {
}
template <typename T>
__device__ size_t wait_until_any_vector(T* ptr, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T* vals){
return 0;
}
template <typename T>
__device__ size_t wait_until_some_vector(T* ptr, size_t nelems,
size_t* indices,
const int *status,
roc_shmem_cmps cmp, T* vals) {
}
template <typename T>
__device__ int test(T* ptr, roc_shmem_cmps cmp, T val) {
return 0;
}
} // namespace rocshmem
#endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_
projects/rocshmem/src/ipc/ipc_context_proxy.hpp
+3 -2
مشاهده پرونده
@@ -41,10 +41,11 @@ class IPCDefaultContextProxy {
/*
* Placement new the memory which is allocated by proxy_
*/
explicit IPCDefaultContextProxy(IPCBackend* backend) : constructed_{true} {
explicit IPCDefaultContextProxy(IPCBackend* backend, TeamInfo *tinfo)
: constructed_{true} {
auto ctx{proxy_.get()};
new (ctx) IPCContext(reinterpret_cast<Backend*>(backend));
roc_shmem_ctx_t local{ctx, nullptr};
roc_shmem_ctx_t local{ctx, tinfo};
set_internal_ctx(&local);
}
projects/rocshmem/src/ipc_policy.hpp
+12 -12
مشاهده پرونده
@@ -64,38 +64,38 @@ class IpcOnImpl {
__device__ void ipcCopy_wave(void *dst, void *src, size_t size);
__device__ void ipcFence() { __threadfence(); }
__device__ void ipcFence() { __threadfence_system(); }
template <typename T>
__device__ T ipcAMOFetchAdd(T *val, T value) {
return __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
return __hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST,
__HIP_MEMORY_SCOPE_SYSTEM);
}
template <typename T>
__device__ T ipcAMOFetchCas(T *val, T cond, T value) {
__hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED,
__ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
__hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST,
__HIP_MEMORY_SCOPE_SYSTEM);
return cond;
}
template <typename T>
__device__ void ipcAMOAdd(T *val, T value) {
__hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
__hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST,
__HIP_MEMORY_SCOPE_SYSTEM);
}
template <typename T>
__device__ void ipcAMOCas(T *val, T cond, T value) {
__hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED,
__ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_AGENT);
__hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST,
__HIP_MEMORY_SCOPE_SYSTEM);
}
template <typename T>
__device__ void ipcAMOSet(T *val, T value) {
__hip_atomic_store(val, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
__hip_atomic_store(val, value, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SYSTEM);
}
__device__ void zero_byte_read(int pe) {
projects/rocshmem/src/memory/heap_type.hpp
+1 -1
مشاهده پرونده
@@ -40,7 +40,7 @@ namespace rocshmem {
#if defined USE_MANAGED_HEAP
using HEAP_T = HeapMemory<HIPAllocatorManaged>;
#elif defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
#elif defined USE_COHERENT_HEAP
using HEAP_T = HeapMemory<HIPAllocator>;
#elif defined USE_HOST_HEAP
using HEAP_T = HeapMemory<HostAllocator>;
projects/rocshmem/src/memory/notifier.hpp
+67 -30
مشاهده پرونده
@@ -20,59 +20,96 @@
* IN THE SOFTWARE.
*****************************************************************************/
/**
* @file notifier.hpp
*
* @brief Contains the notification memory space for threads to communicate
* results with one another.
*
* Assume one thread does work on behalf of other threads (as a leader) and
* that work needs to be communicated to the other threads as a result.
* To expose the result, the threads need to share a memory space where the
* result can be written. The leader thread writes the result out to this
* memory space and all threads synchronize on it.
*
* This class allows the leader thread to notify other threads of the update.
*/
#ifndef LIBRARY_SRC_MEMORY_NOTIFIER_HPP_
#define LIBRARY_SRC_MEMORY_NOTIFIER_HPP_
#include "../device_proxy.hpp"
#include "../util.hpp"
#include "../atomic.hpp"
namespace rocshmem {
template<detail::atomic::rocshmem_memory_scope scope>
class Notifier {
public:
__device__ uint64_t read() { return value_; }
__device__ void write(uint64_t val) {
if (is_thread_zero_in_block()) {
value_ = val;
}
publish();
public:
__device__ uint64_t load() {
return detail::atomic::load<uint64_t, scope>(&value_, orders_);
}
__device__ void done() { __syncthreads(); }
__device__ void store(uint64_t val) {
detail::atomic::store<uint64_t, scope>(&value_, val, orders_);
}
private:
__device__ void publish() {
if (is_thread_zero_in_block()) {
__threadfence();
__device__ void fence() {
detail::atomic::threadfence<scope>();
}
__device__ void sync() {
if constexpr (scope == detail::atomic::memory_scope_thread ||
scope == detail::atomic::memory_scope_wavefront) {
return;
}
if constexpr (scope == detail::atomic::memory_scope_workgroup) {
__syncthreads();
return;
}
if constexpr (scope == detail::atomic::memory_scope_system) {
assert(false);
return;
}
uint32_t done {signal_ + 1};
__syncthreads();
uint32_t retval {0};
bool executor {!threadIdx.x && !threadIdx.y && !threadIdx.z};
if (executor) {
retval = detail::atomic::fetch_add<uint32_t, uint32_t, scope>(&count_, 1, orders_);
fence();
}
__syncthreads();
if (retval == ((gridDim.x * gridDim.y * gridDim.z) - 1)) {
if (executor) {
detail::atomic::store<uint32_t, scope>(&count_, 0, orders_);
fence();
detail::atomic::fetch_add<uint32_t, uint32_t, scope>(&signal_, 1, orders_);
}
}
if (executor) {
while (detail::atomic::load<uint32_t, scope>(&signal_, orders_) != done) {
;
}
}
__syncthreads();
}
private:
detail::atomic::rocshmem_memory_orders orders_{};
uint64_t value_{};
uint32_t signal_ {};
uint32_t count_ {};
};
template <typename ALLOCATOR>
template <typename ALLOCATOR, detail::atomic::rocshmem_memory_scope scope>
class NotifierProxy {
using ProxyT = DeviceProxy<ALLOCATOR, Notifier, 1>;
using ProxyT = DeviceProxy<ALLOCATOR, Notifier<scope>>;
public:
__host__ __device__ Notifier* get() { return proxy_.get(); }
NotifierProxy() {
new (proxy_.get()) Notifier<scope>();
}
~NotifierProxy() {
proxy_.get()->~Notifier<scope>();
}
__host__ __device__ Notifier<scope>* get() { return proxy_.get(); }
private:
ProxyT proxy_{};
projects/rocshmem/src/memory/slab_heap.cpp
+6 -3
مشاهده پرونده
@@ -75,9 +75,12 @@ __device__ void SlabHeap::malloc(void** ptr, size_t size) {
* Notify other threads in block about the allocation result.
*/
auto notifier{notifier_.get()};
notifier->write(ptr_deref_u64);
uint64_t notification_u64{notifier->read()};
notifier->done();
if (!threadIdx.x) {
notifier->store(ptr_deref_u64);
notifier->fence();
}
__syncthreads();
uint64_t notification_u64{notifier->load()};
/*
* Write to the ptr parameter (to return it back up the call stack).
projects/rocshmem/src/memory/slab_heap.hpp
+1 -1
مشاهده پرونده
@@ -48,7 +48,7 @@ class SlabHeap {
/**
* @brief Helper type for notifier
*/
using NOTIFIER_PROXY_T = NotifierProxy<HIPAllocator>;
using NOTIFIER_PROXY_T = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_workgroup>;
/**
* @brief Helper type for notifier
projects/rocshmem/src/reverse_offload/backend_ro.cpp
+3 -1
مشاهده پرونده
@@ -94,7 +94,9 @@ ROBackend::ROBackend(MPI_Comm comm)
default_block_handle_proxy_ = DefaultBlockHandleProxyT(
bp->g_ret, bp->atomic_ret, &queue_, &ipcImpl, hdp_proxy_.get());
default_context_proxy_ = DefaultContextProxyT(this);
TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world;
default_context_proxy_ = DefaultContextProxyT(this, tinfo);
block_handle_proxy_ = BlockHandleProxyT(bp->g_ret, bp->atomic_ret, &queue_,
&ipcImpl, hdp_proxy_.get());
projects/rocshmem/src/reverse_offload/context_proxy.hpp
+3 -2
مشاهده پرونده
@@ -42,10 +42,11 @@ class DefaultContextProxy {
/*
* Placement new the memory which is allocated by proxy_
*/
explicit DefaultContextProxy(ROBackend* backend) : constructed_{true} {
explicit DefaultContextProxy(ROBackend* backend, TeamInfo *tinfo)
: constructed_{true} {
auto ctx{proxy_.get()};
new (ctx) ROContext(reinterpret_cast<Backend*>(backend), -1);
roc_shmem_ctx_t local{ctx, nullptr};
roc_shmem_ctx_t local{ctx, tinfo};
set_internal_ctx(&local);
}
projects/rocshmem/src/reverse_offload/context_ro_device.cpp
+3 -18
مشاهده پرونده
@@ -50,6 +50,9 @@ __host__ ROContext::ROContext(Backend *b, size_t block_id)
block_handle = &block_base[block_id];
}
ro_net_win_id = block_id % backend->ro_window_proxy_->MAX_NUM_WINDOWS;
ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
ipcImpl_.shm_size = b->ipcImpl.shm_size;
}
__device__ void ROContext::putmem(void *dest, const void *source, size_t nelems,
@@ -389,24 +392,6 @@ __device__ bool enough_space(BlockHandle *h, uint64_t required) {
return (h->queue_size - (h->write_index - h->read_index)) >= required;
}
__device__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value,
volatile uint64_t *read_value) {
__asm__ volatile(
"global_load_dwordx2 %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
}
__device__ void refresh_volatile_sbyte(volatile int *assigned_value,
volatile char *read_value) {
__asm__ volatile(
"global_load_sbyte %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)"
: "=v"(*assigned_value)
: "v"(read_value));
}
__device__ void acquire_lock(BlockHandle *handle) {
while(atomicCAS((uint64_t *)&handle->lock, 0, 1) == 1) ;
}
projects/rocshmem/src/roc_shmem_gpu.cpp
+6 -4
مشاهده پرونده
@@ -262,9 +262,9 @@ __device__ int roc_shmem_wg_ctx_create(long option, roc_shmem_ctx_t *ctx) {
GPU_DPRINTF("Function: roc_shmem_ctx_create\n");
bool result{true};
if (get_flat_block_id() == 0) {
ctx->team_opaque = reinterpret_cast<TeamInfo *>(ROC_SHMEM_CTX_DEFAULT.team_opaque);
device_backend_proxy->create_ctx(option, ctx);
reinterpret_cast<Context *>(ctx->ctx_opaque)->setFence(option);
ctx->team_opaque = nullptr;
}
__syncthreads();
return result == true ? 0 : -1;
@@ -279,11 +279,11 @@ __device__ int roc_shmem_wg_team_create_ctx(roc_shmem_team_t team, long options,
bool result{true};
if (get_flat_block_id() == 0) {
result = device_backend_proxy->create_ctx(options, ctx);
reinterpret_cast<Context *>(ctx->ctx_opaque)->setFence(options);
Team *team_obj{get_internal_team(team)};
TeamInfo *info_wrt_world = team_obj->tinfo_wrt_world;
ctx->team_opaque = info_wrt_world;
result = device_backend_proxy->create_ctx(options, ctx);
reinterpret_cast<Context *>(ctx->ctx_opaque)->setFence(options);
}
__syncthreads();
@@ -412,7 +412,9 @@ __device__ void roc_shmem_ctx_fence(roc_shmem_ctx_t ctx) {
__device__ void roc_shmem_ctx_fence(roc_shmem_ctx_t ctx, int pe) {
GPU_DPRINTF("Function: roc_shmem_ctx_fence\n");
get_internal_ctx(ctx)->fence(pe);
int pe_in_world = translate_pe(ctx, pe);
get_internal_ctx(ctx)->fence(pe_in_world);
}
__device__ void roc_shmem_ctx_quiet(roc_shmem_ctx_t ctx) {
projects/rocshmem/src/util.cpp
-9
مشاهده پرونده
@@ -40,15 +40,6 @@ typedef struct device_agent {
std::vector<device_agent_t> gpu_agents;
std::vector<device_agent_t> cpu_agents;
__device__ uint64_t __read_clock() {
uint64_t clock{};
asm volatile(
"s_memrealtime %0\n"
"s_waitcnt lgkmcnt(0)\n"
: "=s"(clock));
return clock;
}
hsa_status_t rocm_hsa_amd_memory_pool_callback(
hsa_amd_memory_pool_t memory_pool, void* data) {
hsa_amd_memory_pool_global_flag_t pool_flag{};
projects/rocshmem/src/util.hpp
+10 -88
مشاهده پرونده
@@ -29,60 +29,12 @@
#include <cstdio>
#include "assembly.hpp"
#include "config.h" // NOLINT(build/include_subdir)
#include "constants.hpp"
namespace rocshmem {
#define DO_PRAGMA(x) _Pragma(#x)
#define NOWARN(warnoption, ...) \
DO_PRAGMA(GCC diagnostic push) \
DO_PRAGMA(GCC diagnostic ignored #warnoption) \
__VA_ARGS__ \
DO_PRAGMA(GCC diagnostic pop)
__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
int ret;
__asm__ volatile(
"global_load_ubyte %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
return ret;
}
/* Ignore the warning about deprecated volatile.
* The only usage of volatile is to force the compiler to generate
* the assembly instruction. If volatile is omitted, the compiler
* will NOT generate the non-temporal load or the waitcnt.
*/
// clang-format off
NOWARN(-Wdeprecated-volatile,
template <typename T> __device__ __forceinline__ T uncached_load(T* src) {
T ret;
switch (sizeof(T)) {
case 4:
__asm__ volatile(
"global_load_dword %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
break;
case 8:
__asm__ volatile(
"global_load_dwordx2 %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)"
: "=v"(ret)
: "v"(src));
break;
default:
break;
}
return ret;
}
)
// clang-format on
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST)
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST)
@@ -96,8 +48,6 @@ NOWARN(-Wdeprecated-volatile,
} \
}
#define SFENCE() asm volatile("sfence" ::: "memory")
#ifdef DEBUG
#define DPRINTF(...) \
do { \
@@ -123,19 +73,6 @@ NOWARN(-Wdeprecated-volatile,
extern const int gpu_clock_freq_mhz;
/* Device-side internal functions */
__device__ __forceinline__ void __roc_inv() { asm volatile("buffer_wbinvl1;"); }
__device__ __forceinline__ void __roc_flush() {
#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP
#if __gfx90a__
asm volatile("s_dcache_wb;");
asm volatile("buffer_wbl2;");
#endif
#endif
}
__device__ uint64_t __read_clock();
__device__ __forceinline__ uint32_t lowerID() {
return __ffsll(__ballot(1)) - 1;
}
@@ -156,6 +93,7 @@ __device__ __forceinline__ bool is_thread_zero_in_block() {
__device__ __forceinline__ bool is_block_zero_in_grid() {
return hipBlockIdx_x == 0 && hipBlockIdx_y == 0 && hipBlockIdx_z == 0;
}
/*
* Returns the number of threads in the caller's flattened thread block.
*/
@@ -163,6 +101,13 @@ __device__ __forceinline__ int get_flat_block_size() {
return hipBlockDim_x * hipBlockDim_y * hipBlockDim_z;
}
/*
* Returns the number of threads in the caller's flattened grid.
*/
__device__ __forceinline__ int get_flat_grid_size() {
return get_flat_block_size() * hipGridDim_x * hipGridDim_y * hipGridDim_z;
}
/*
* Returns the flattened thread index of the calling thread within its
* thread block.
@@ -212,7 +157,7 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) {
while (atomicCAS(print_lock, 0, 1) == 1) {
}
printf("WG (%lu, %lu, %lu) TH (%lu, %lu, %lu) ", hipBlockIdx_x,
printf("WG (%u, %u, %u) TH (%u, %u, %u) ", hipBlockIdx_x,
hipBlockIdx_y, hipBlockIdx_z, hipThreadIdx_x, hipThreadIdx_y,
hipThreadIdx_z);
printf(fmt, args...);
@@ -222,29 +167,6 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) {
}
}
__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
int size) {
switch (size) {
case 2: {
int16_t val16{*(reinterpret_cast<int16_t*>(val))};
asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
break;
}
case 4: {
int32_t val32{*(reinterpret_cast<int32_t*>(val))};
asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
break;
}
case 8: {
int64_t val64{*(reinterpret_cast<int64_t*>(val))};
asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
break;
}
default:
break;
}
}
__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)};
projects/rocshmem/tests/functional_tests/sync_tester.cpp
+12
مشاهده پرونده
@@ -83,9 +83,21 @@ void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
roc_shmem_team_split_strided(ROC_SHMEM_TEAM_WORLD, 0, 1, n_pes, nullptr, 0,
&team_sync_world_dup);
#ifdef USE_COOPERATIVE_GROUPS
void* kernelParams[] = {(void*)&loop,
(void*)&args.skip,
(void*)&timer,
(void*)&_type,
(void*)&_shmem_context,
(void*)&team_sync_world_dup};
CHECK_HIP(hipLaunchCooperativeKernel(SyncTest, gridSize, blockSize,
kernelParams, shared_bytes, stream));
#else
hipLaunchKernelGGL(SyncTest, gridSize, blockSize, shared_bytes, stream, loop,
args.skip, timer, _type, _shmem_context,
team_sync_world_dup);
#endif /* USE_COOPERATIVE_GROUPS */
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
projects/rocshmem/tests/unit_tests/CMakeLists.txt
+1
مشاهده پرونده
@@ -90,6 +90,7 @@ target_sources(
free_list_gtest.cpp
#context_ipc_gtest.cpp
ipc_impl_simple_coarse_gtest.cpp
ipc_impl_simple_fine_gtest.cpp
)
###############################################################################
projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp
+1
مشاهده پرونده
@@ -33,6 +33,7 @@ TEST_F(IPCImplSimpleCoarseTestFixture, MPI_num_pes) {
}
TEST_F(IPCImplSimpleCoarseTestFixture, IPC_bases) {
ASSERT_NE(ipc_impl_.ipc_bases, nullptr);
for(int i{0}; i < mpi_.num_pes(); i++) {
ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr);
}
projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp
-2
مشاهده پرونده
@@ -218,8 +218,6 @@ class IPCImplSimpleCoarseTestFixture : public ::testing::Test {
protected:
std::vector<int> golden_;
std::vector<int> output_;
HEAP_T heap_mem_ {};
MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()};
projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp
+1037
مشاهده پرونده
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است Diff را بارگزاری کن
projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp
+344
مشاهده پرونده
@@ -0,0 +1,344 @@
/******************************************************************************
* Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
*
* 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_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
#define ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
#include "gtest/gtest.h"
#include <numeric>
#include <mpi.h>
#include "../src/atomic.hpp"
#include "../src/ipc_policy.hpp"
#include "../src/memory/notifier.hpp"
#include "../src/memory/symmetric_heap.hpp"
#include "../src/util.hpp"
namespace rocshmem {
enum TestType {
READ = 0,
WRITE = 1
};
const uint32_t SIGNAL_OFFSET {67108864};
__device__
void
validator(bool *error, int *golden, int *dest, size_t bytes) {
size_t elements {bytes / sizeof(int)};
for (int i {get_flat_id()}; i < elements; i += get_flat_grid_size()) {
if (golden[i] != dest[i]) {
printf("golden[%d] %d != dest[%d] %d\n", i, golden[i], i, dest[i]);
*error = true;
}
}
}
template <typename NotifierT>
__global__
void
kernel_put_with_signal_validator(bool *error, int *golden, int *dest, size_t bytes, NotifierT *notifier) {
detail::atomic::rocshmem_memory_orders orders{};
if (!get_flat_id()) {
while (detail::atomic::load<int, detail::atomic::memory_scope_system>(dest + SIGNAL_OFFSET, orders) == 0) {
;
}
}
notifier->sync();
validator(error, golden, dest, bytes);
}
template <typename NotifierT>
__global__
void
kernel_simple_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
if (!get_flat_id()) {
ipc_impl->ipcCopy(dest, src, bytes);
ipc_impl->ipcFence();
if (test == WRITE) {
ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
}
}
if (test == READ) {
notifier->sync();
validator(error, golden, dest, bytes);
}
}
template <typename NotifierT>
__global__
void
kernel_simple_fine_copy_wg(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
if (!blockIdx.x) {
ipc_impl->ipcCopy_wg(dest, src, bytes);
ipc_impl->ipcFence();
if (test == WRITE) {
if (!threadIdx.x) {
ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
}
}
}
if (test == READ) {
notifier->sync();
validator(error, golden, dest, bytes);
}
}
template <typename NotifierT>
__global__
void
kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) {
if (!blockIdx.x && threadIdx.x < 64) {
ipc_impl->ipcCopy_wave(dest, src, bytes);
ipc_impl->ipcFence();
if (test == WRITE) {
if (!threadIdx.x) {
ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1);
}
}
}
__syncthreads();
if (test == READ) {
notifier->sync();
validator(error, golden, dest, bytes);
}
}
class IPCImplSimpleFineTestFixture : public ::testing::Test {
using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
using FN_T1 = void (*)(IpcImpl*, bool*, int*, int*, int*, size_t, TestType, NotifierT*);
using FN_T2 = void (*)(bool*, int*, int*, size_t, NotifierT*);
public:
IPCImplSimpleFineTestFixture() {
ipc_impl_.ipcHostInit(mpi_.my_pe(), mpi_.get_heap_bases() , MPI_COMM_WORLD);
assert(ipc_impl_dptr_ == nullptr);
hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl));
CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_, sizeof(IpcImpl), hipMemcpyHostToDevice));
assert(error_dptr_ == nullptr);
hip_allocator_.allocate((void**)&error_dptr_, sizeof(bool));
*error_dptr_ = false;
}
~IPCImplSimpleFineTestFixture() {
if (ipc_impl_dptr_) {
hip_allocator_.deallocate(ipc_impl_dptr_);
}
if (error_dptr_) {
hip_allocator_.deallocate(error_dptr_);
}
if (golden_dptr_) {
hip_allocator_.deallocate(golden_dptr_);
}
ipc_impl_.ipcHostStop();
}
void launch(FN_T1 f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes, TestType test) {
f<<<grid, block>>>(ipc_impl_dptr_, error_dptr_, golden_dptr_, src, dest, bytes, test, notifier_.get());
CHECK_HIP(hipStreamSynchronize(nullptr));
}
void launch(FN_T2 f, const dim3 grid, const dim3 block, int* dest, size_t bytes) {
f<<<grid, block>>>(error_dptr_, golden_dptr_, dest, bytes, notifier_.get());
CHECK_HIP(hipStreamSynchronize(nullptr));
}
void write(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
initialize_src_buffer(WRITE);
copy(WRITE, grid, block);
check_device_validation_errors(WRITE);
}
void write_wg(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
initialize_src_buffer(WRITE);
copy_wg(WRITE, grid, block);
check_device_validation_errors(WRITE);
}
void write_wave(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(WRITE);
initialize_src_buffer(WRITE);
copy_wave(WRITE, grid, block);
check_device_validation_errors(WRITE);
}
void read(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
initialize_src_buffer(READ);
copy(READ, grid, block);
check_device_validation_errors(READ);
}
void read_wg(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
initialize_src_buffer(READ);
copy_wg(READ, grid, block);
check_device_validation_errors(READ);
}
void read_wave(const dim3 grid, const dim3 block, size_t elems) {
iota_golden(elems);
initialize_signal(READ);
initialize_src_buffer(READ);
copy_wave(READ, grid, block);
check_device_validation_errors(READ);
}
void iota_golden(size_t elems) {
golden_.resize(elems);
std::iota(golden_.begin(), golden_.end(), 0);
assert(golden_dptr_ == nullptr);
size_t golden_dptr_bytes {golden_.size() * sizeof(int)};
hip_allocator_.allocate((void**)&golden_dptr_, golden_dptr_bytes);
CHECK_HIP(hipMemcpy(golden_dptr_, golden_.data(), golden_dptr_bytes, hipMemcpyHostToDevice));
}
void validate_golden(size_t elems) {
ASSERT_EQ(golden_.size(), elems);
for (int i{0}; i < golden_.size(); i++) {
ASSERT_EQ(golden_[i], i);
}
}
void initialize_signal(TestType test) {
bool is_write_test = test;
if (is_write_test && mpi_.my_pe() == 0) {
int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
*(dest + SIGNAL_OFFSET) = 0;
}
}
void initialize_src_buffer(TestType test) {
if (!pe_initializes_src_buffer(test)) {
return;
}
size_t bytes = golden_.size() * sizeof(int);
auto dev_src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_.my_pe()]);
CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice));
}
bool pe_initializes_src_buffer(TestType test) {
bool is_write_test = test;
bool is_read_test = !test;
return (is_write_test && mpi_.my_pe() == 0) ||
(is_read_test && mpi_.my_pe() == 1);
}
void execute(TestType test, FN_T1 fn, const dim3 grid, const dim3 block) {
size_t bytes = golden_.size() * sizeof(int);
if (mpi_.my_pe()) {
mpi_.barrier();
if (test == WRITE) {
int *dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
FN_T2 val_fn = kernel_put_with_signal_validator;
launch(val_fn, grid, block, dest, bytes);
}
mpi_.barrier();
return;
}
int *src{nullptr};
int *dest{nullptr};
if (test == WRITE) {
src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
} else {
src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[0]);
}
mpi_.barrier();
launch(fn, grid, block, src, dest, bytes, test);
mpi_.barrier();
}
void copy(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy, grid, block);
}
void copy_wg(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy_wg, grid, block);
}
void copy_wave(TestType test, dim3 grid, dim3 block) {
execute(test, kernel_simple_fine_copy_wave, grid, block);
}
void check_device_validation_errors(TestType test) {
if (!pe_validates_dest_buffer(test)) {
return;
}
ASSERT_EQ(*error_dptr_, false);
}
void validate_dest_buffer(TestType test) {
if (!pe_validates_dest_buffer(test)) {
return;
}
auto dev_dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[mpi_.my_pe()]);
for (int i{0}; i < golden_.size(); i++) {
ASSERT_EQ(golden_[i], dev_dest[i]);
}
}
bool pe_validates_dest_buffer(TestType test) {
return !pe_initializes_src_buffer(test);
}
protected:
HIPDefaultFinegrainedAllocator hip_allocator_ {};
NotifierProxyT notifier_ {};
HEAP_T heap_mem_ {};
MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()};
std::vector<int> golden_;
int *golden_dptr_ {nullptr};
IpcImpl ipc_impl_ {};
IpcImpl *ipc_impl_dptr_ {nullptr};
bool *error_dptr_ {nullptr};
};
} // namespace rocshmem
#endif // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP
projects/rocshmem/tests/unit_tests/notifier_gtest.cpp
+75 -7
مشاهده پرونده
@@ -28,30 +28,98 @@ using namespace rocshmem;
******************************* Fixture Tests *******************************
*****************************************************************************/
TEST_F(NotifierTestFixture, run_all_threads_once_1_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_1_1) {
run_all_threads_once(1, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_2_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_2_1) {
run_all_threads_once(2, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_64_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_64_1) {
run_all_threads_once(64, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_128_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_128_1) {
run_all_threads_once(128, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_256_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_256_1) {
run_all_threads_once(256, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_512_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_512_1) {
run_all_threads_once(512, 1);
}
TEST_F(NotifierTestFixture, run_all_threads_once_1024_1) {
TEST_F(NotifierBlockTestFixture, run_all_threads_once_1024_1) {
run_all_threads_once(1024, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_1) {
run_all_threads_once(1, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_2_1) {
run_all_threads_once(2, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_64_1) {
run_all_threads_once(64, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_128_1) {
run_all_threads_once(128, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_256_1) {
run_all_threads_once(256, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_512_1) {
run_all_threads_once(512, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_1) {
run_all_threads_once(1024, 1);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_2) {
run_all_threads_once(1, 2);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_2) {
run_all_threads_once(1024, 2);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_4) {
run_all_threads_once(1, 4);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_4) {
run_all_threads_once(1024, 4);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_8) {
run_all_threads_once(1, 8);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_8) {
run_all_threads_once(1024, 8);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_32) {
run_all_threads_once(1, 32);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_32) {
run_all_threads_once(1024, 32);
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_38) {
run_all_threads_once(1, 38); // MI300 CPX
}
TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_38) {
run_all_threads_once(1024, 38); // MI300 CPX
}
projects/rocshmem/tests/unit_tests/notifier_gtest.hpp
+52 -34
مشاهده پرونده
@@ -44,66 +44,45 @@ static const uint64_t NOTIFIER_OFFSET {0x100B00};
inline __device__
void
write_to_memory(uint8_t* raw_memory) {
auto thread_idx {get_flat_block_id()};
auto thread_idx {get_flat_id()};
raw_memory[thread_idx] = THREAD_VALUE;
__threadfence();
}
template <typename NotifierT>
__global__
void
all_threads_once(uint8_t* raw_memory,
Notifier* notifier) {
notifier->write(NOTIFIER_OFFSET);
uint64_t offset_u64 {notifier->read()};
notifier->done();
NotifierT * notifier) {
if (!get_flat_id()) {
notifier->store(NOTIFIER_OFFSET);
notifier->fence();
}
notifier->sync();
uint64_t offset_u64 {notifier->load()};
uint64_t raw_memory_u64 {reinterpret_cast<uint64_t>(raw_memory)};
uint64_t address_u64 {raw_memory_u64 + offset_u64};
uint8_t* address {reinterpret_cast<uint8_t*>(address_u64)};
write_to_memory(address);
__syncthreads();
}
class NotifierTestFixture : public ::testing::Test {
using NotifierProxyT = NotifierProxy<HIPAllocator>;
class NotifierBase : public ::testing::Test {
public:
NotifierTestFixture() {
NotifierBase() {
assert(raw_memory_ == nullptr);
hip_allocator_.allocate((void**)&raw_memory_, GIBIBYTE_);
assert(raw_memory_);
}
~NotifierTestFixture() {
~NotifierBase() {
if (raw_memory_) {
hip_allocator_.deallocate(raw_memory_);
}
}
void
run_all_threads_once(uint32_t x_block_dim,
uint32_t x_grid_dim) {
const dim3 hip_blocksize(x_block_dim, 1, 1);
const dim3 hip_gridsize(x_grid_dim, 1, 1);
hipLaunchKernelGGL(all_threads_once,
hip_gridsize,
hip_blocksize,
0,
nullptr,
raw_memory_,
notifier_.get());
hipError_t return_code = hipStreamSynchronize(nullptr);
if (return_code != hipSuccess) {
printf("Failed in stream synchronize\n");
assert(return_code == hipSuccess);
}
size_t number_threads {x_block_dim * x_grid_dim};
verify(size_t number_threads) {
uint8_t* offset_addr {compute_offset_addr()};
for (size_t i {0}; i < number_threads; i++) {
ASSERT_EQ(offset_addr[i], THREAD_VALUE);
}
@@ -136,12 +115,51 @@ class NotifierTestFixture : public ::testing::Test {
*/
uint8_t *raw_memory_ {nullptr};
};
class NotifierBlockTestFixture : public NotifierBase {
using NotifierT = Notifier<detail::atomic::memory_scope_workgroup>;
using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_workgroup>;
public:
void
run_all_threads_once(uint32_t x_block_dim,
uint32_t x_grid_dim) {
new (notifier_.get()) NotifierT();
const dim3 block(x_block_dim, 1, 1);
const dim3 grid(x_grid_dim, 1, 1);
all_threads_once<NotifierT><<<grid, block>>>(raw_memory_, notifier_.get());
CHECK_HIP(hipStreamSynchronize(nullptr));
verify(x_block_dim * x_grid_dim);
}
/**
* @brief Used to broadcast base offset for writing.
*/
NotifierProxyT notifier_ {};
};
class NotifierAgentTestFixture : public NotifierBase {
using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
public:
void
run_all_threads_once(uint32_t x_block_dim,
uint32_t x_grid_dim) {
new (notifier_.get()) NotifierT();
const dim3 block(x_block_dim, 1, 1);
const dim3 grid(x_grid_dim, 1, 1);
all_threads_once<NotifierT><<<grid, block>>>(raw_memory_, notifier_.get());
CHECK_HIP(hipStreamSynchronize(nullptr));
verify(x_block_dim * x_grid_dim);
}
/**
* @brief Used to broadcast base offset for writing.
*/
NotifierProxyT notifier_ {};
};
} // namespace rocshmem
projects/rocshmem/tests/unit_tests/shmem_gtest.cpp
+20 -1
مشاهده پرونده
@@ -21,8 +21,27 @@
*****************************************************************************/
#include "gtest/gtest.h"
#include <mpi.h>
int main(int argc, char **argv) {
::testing::InitGoogleTest(&argc, argv);
return RUN_ALL_TESTS();
int initialized;
MPI_Initialized(&initialized);
if (!initialized) {
int provided;
MPI_Init_thread(nullptr, nullptr, MPI_THREAD_MULTIPLE, &provided);
if (provided != MPI_THREAD_MULTIPLE) {
std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n";
}
}
int ret_val = RUN_ALL_TESTS();
int finalized{0};
MPI_Finalized(&finalized);
if (!finalized) {
MPI_Finalize();
}
return ret_val;
}