Update atomic functional tests (#262)
* feat: implement function to return number of blocks in grid.
* test: update atomics functional tests
- Standard atomic tests: `atomic_add`, `atomic_inc`, `fetch_atomic_add`, `fetch_atomic_inc`, and `fetch_compare_and_swap`
- Bitwise atomic tests: `atomic_and`, `atomic_or`, `atomic_xor`, fetch_atomic_and`, `fetch_atomic_or`, and `fetch_atomic_xor`
- Extended atomic tests: `atomic_fetch`, `atomic_set`, and `atomic_swap`
* Added two different address modes for atomics.
* Added all supported data types for atomics tests.
[ROCm/rocshmem commit: 0a4f8a83b9]
This commit is contained in:
Avinash Kethineedi
committed by
GitHub
GitHub
parent
ac13b22edc
commit
e31b4d42e5
@@ -32,33 +32,54 @@ using namespace rocshmem;
|
||||
/* Declare the global kernel template with a generic implementation */
|
||||
template <typename T>
|
||||
__global__ void AMOBitwiseTest(int loop, int skip, long long int *start_time,
|
||||
long long int *end_time, char *r_buf,
|
||||
T *s_buf, T *ret_val, TestType type,
|
||||
long long int *end_time, T *dest, T *ret_val,
|
||||
AddrMode addr_mode, TestType type,
|
||||
ShmemContextType ctx_type) {
|
||||
return;
|
||||
}
|
||||
|
||||
template <class T>
|
||||
__device__ inline T* compute_target_ptr(T* base_ptr, AddrMode addr_mode,
|
||||
int wg_idx, int itr, int n_wgs) {
|
||||
// PerBlock: element = wg_idx, with n_wgs elements per loop
|
||||
// PerGrid : single element shared by the whole grid per loop
|
||||
if (addr_mode == AddrMode::PerBlock) {
|
||||
size_t offset = wg_idx + itr * n_wgs;
|
||||
return base_ptr + offset;
|
||||
} else { // PerGrid
|
||||
return base_ptr + itr;
|
||||
}
|
||||
}
|
||||
|
||||
/******************************************************************************
|
||||
* HOST TESTER CLASS METHODS
|
||||
*****************************************************************************/
|
||||
template <typename T>
|
||||
AMOBitwiseTester<T>::AMOBitwiseTester(TesterArguments args) : Tester(args) {
|
||||
CHECK_HIP(hipMalloc((void **)&_ret_val, args.max_msg_size * args.num_wgs));
|
||||
_r_buf = (char *)rocshmem_malloc(args.max_msg_size);
|
||||
_s_buf = (T *)rocshmem_malloc(args.max_msg_size * args.num_wgs);
|
||||
n_out = (args.addr_mode == AddrMode::PerBlock) ? args.num_wgs : 1;
|
||||
n_in = args.num_wgs * args.wg_size;
|
||||
n_loops = args.loop + args.skip;
|
||||
|
||||
// One return per *thread* per loop
|
||||
CHECK_HIP(hipMalloc((void **)&ret_val, args.max_msg_size * n_in * n_loops));
|
||||
|
||||
dest = (T *)rocshmem_malloc(args.max_msg_size * n_out * n_loops);
|
||||
if (dest == nullptr) {
|
||||
std::cerr << "Error allocating memory from symmetric heap" << std::endl;
|
||||
std::cerr << "dest: " << (void*)dest << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
AMOBitwiseTester<T>::~AMOBitwiseTester() {
|
||||
rocshmem_free(_r_buf);
|
||||
CHECK_HIP(hipFree(_ret_val));
|
||||
CHECK_HIP(hipFree(ret_val));
|
||||
rocshmem_free(dest);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void AMOBitwiseTester<T>::resetBuffers(size_t size) {
|
||||
memset(_r_buf, 0, args.max_msg_size);
|
||||
memset(_ret_val, 0, args.max_msg_size * args.num_wgs);
|
||||
memset(_s_buf, 0, args.max_msg_size * args.num_wgs);
|
||||
memset(ret_val, 0, args.max_msg_size * n_in * n_loops);
|
||||
memset(dest, 0, args.max_msg_size * n_out * n_loops);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
@@ -67,59 +88,158 @@ void AMOBitwiseTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
|
||||
size_t shared_bytes = 0;
|
||||
|
||||
hipLaunchKernelGGL(AMOBitwiseTest, gridsize, blocksize, shared_bytes, stream,
|
||||
loop, args.skip, start_time, end_time, _r_buf, _s_buf,
|
||||
_ret_val, _type, _shmem_context);
|
||||
args.loop, args.skip, start_time, end_time, dest,
|
||||
ret_val, args.addr_mode, _type, _shmem_context);
|
||||
|
||||
_gridSize = gridsize;
|
||||
num_msgs = (loop + args.skip) * gridsize.x;
|
||||
num_timed_msgs = loop;
|
||||
num_msgs = n_loops * gridsize.x * blocksize.x;
|
||||
num_timed_msgs = args.loop * gridsize.x * blocksize.x;
|
||||
}
|
||||
|
||||
template <typename G>
|
||||
void fail_eq(const G& got, const G& exp) {
|
||||
std::cerr << "data validation error\n"
|
||||
<< "got " << got << ", expected " << exp << std::endl;
|
||||
std::exit(-1);
|
||||
}
|
||||
|
||||
// Map (loop, elem_idx) -> dest[] index for current address mode.
|
||||
template <typename T>
|
||||
int AMOBitwiseTester<T>::destIndex(int l, int elem_idx) const {
|
||||
return (args.addr_mode == AddrMode::PerBlock)
|
||||
? l * static_cast<int>(args.num_wgs) + elem_idx
|
||||
: l; // PerGrid has a single element per loop
|
||||
}
|
||||
|
||||
// Number of output elements to check per loop for current address mode.
|
||||
template <typename T>
|
||||
int AMOBitwiseTester<T>::numElems() const {
|
||||
return (args.addr_mode == AddrMode::PerBlock)
|
||||
? static_cast<int>(args.num_wgs)
|
||||
: 1; // PerGrid
|
||||
}
|
||||
|
||||
// Return pointer to the start of the ret_val “chunk” for (loop, elem_idx)
|
||||
// plus the chunk length for this address mode.
|
||||
template <typename T>
|
||||
std::pair<T*, int> AMOBitwiseTester<T>::retChunk(int l, int elem_idx) const {
|
||||
if (args.addr_mode == AddrMode::PerBlock) {
|
||||
// One chunk per element (workgroup): wg_size returns
|
||||
T* p = ret_val + l * n_in + elem_idx * args.wg_size;
|
||||
int sz = static_cast<int>(args.wg_size);
|
||||
return {p, sz};
|
||||
}
|
||||
// PerGrid: one big chunk per loop (all threads)
|
||||
T* p = ret_val + l * n_in;
|
||||
int sz = static_cast<int>(n_in);
|
||||
return {p, sz};
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void AMOBitwiseTester<T>::verifyDestValues() {
|
||||
const int loops = static_cast<int>(n_loops);
|
||||
const int n_elems = numElems();
|
||||
|
||||
auto check_equal_all = [&](T expected) {
|
||||
for (int l = 0; l < loops; ++l) {
|
||||
for (int elem = 0; elem < n_elems; ++elem) {
|
||||
const int idx = destIndex(l, elem);
|
||||
if (dest[idx] != expected) fail_eq(dest[idx], expected);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// Use all-ones mask for type T
|
||||
const T MASK = static_cast<T>(~T{0});
|
||||
|
||||
switch (_type) {
|
||||
case AMO_AndTestType:
|
||||
case AMO_FetchAndTestType: {
|
||||
// Start at 0; 0 & MASK == 0 regardless of writer count.
|
||||
check_equal_all(static_cast<T>(0));
|
||||
break;
|
||||
}
|
||||
case AMO_OrTestType:
|
||||
case AMO_FetchOrTestType: {
|
||||
// final value is MASK.
|
||||
check_equal_all(MASK);
|
||||
break;
|
||||
}
|
||||
case AMO_XorTestType:
|
||||
case AMO_FetchXorTestType: {
|
||||
// PerBlock: K = wg_size; PerGrid: K = num_wgs * wg_size
|
||||
const int K = (args.addr_mode == AddrMode::PerBlock)
|
||||
? static_cast<int>(args.wg_size)
|
||||
: static_cast<int>(args.num_wgs * args.wg_size);
|
||||
const T expected = (K & 1) ? MASK : static_cast<T>(0);
|
||||
check_equal_all(expected);
|
||||
break;
|
||||
}
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void AMOBitwiseTester<T>::verifyReturnValues() {
|
||||
// Only “fetch-*” types produce return values to validate.
|
||||
if (_type == AMO_AndTestType || _type == AMO_OrTestType ||
|
||||
_type == AMO_XorTestType) return;
|
||||
|
||||
const int loops = static_cast<int>(n_loops);
|
||||
const int n_elems = numElems();
|
||||
const T MASK = static_cast<T>(~T{0});
|
||||
|
||||
for (int l = 0; l < loops; ++l) {
|
||||
for (int elem = 0; elem < n_elems; ++elem) {
|
||||
auto [p, cnt] = retChunk(l, elem);
|
||||
|
||||
// Count distribution of observed old values in this chunk
|
||||
int zeros = 0, masks = 0;
|
||||
for (int i = 0; i < cnt; ++i) {
|
||||
zeros += (p[i] == static_cast<T>(0));
|
||||
masks += (p[i] == MASK);
|
||||
}
|
||||
if (zeros + masks != cnt) {
|
||||
fail_eq(zeros + masks, cnt); // unexpected values present
|
||||
}
|
||||
|
||||
switch (_type) {
|
||||
case AMO_FetchAndTestType:
|
||||
// Old value is 0 (dest stays 0)
|
||||
if (!(zeros == cnt && masks == 0)) fail_eq(zeros, cnt);
|
||||
break;
|
||||
|
||||
case AMO_FetchOrTestType:
|
||||
// Exactly one 0 (the first OR), rest MASK
|
||||
if (!(zeros == 1 && masks == cnt - 1)) fail_eq(zeros, 1);
|
||||
break;
|
||||
|
||||
case AMO_FetchXorTestType: {
|
||||
// returns multiset = { ceil(K/2) zeros, floor(K/2) MASKs }
|
||||
const int exp_zeros = (cnt + 1) / 2; // ceil(cnt/2)
|
||||
const int exp_masks = cnt / 2; // floor(cnt/2)
|
||||
if (!(zeros == exp_zeros && masks == exp_masks)) {
|
||||
fail_eq(zeros, exp_zeros);
|
||||
}
|
||||
// cross-check
|
||||
if ((cnt & 1) && zeros != masks + 1) fail_eq(zeros, masks + 1);
|
||||
if (!(cnt & 1) && zeros != masks) fail_eq(zeros, masks);
|
||||
break;
|
||||
}
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void AMOBitwiseTester<T>::verifyResults(size_t size) {
|
||||
T ret;
|
||||
if (args.myid == 0) {
|
||||
T expected_val = 0;
|
||||
|
||||
switch (_type) {
|
||||
case AMO_FetchAndTestType:
|
||||
expected_val = 0;
|
||||
break;
|
||||
case AMO_AndTestType:
|
||||
expected_val = 0;
|
||||
break;
|
||||
case AMO_FetchOrTestType:
|
||||
expected_val = 0xFFFF;
|
||||
break;
|
||||
case AMO_OrTestType:
|
||||
expected_val = 0xFFFF;
|
||||
break;
|
||||
case AMO_FetchXorTestType:
|
||||
expected_val = 0xFFFF;
|
||||
break;
|
||||
case AMO_XorTestType:
|
||||
expected_val = (num_msgs % 2) ? 0xFFFF : 0;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
|
||||
int fetch_op =
|
||||
(_type == AMO_FetchAndTestType || _type == AMO_FetchOrTestType ||
|
||||
_type == AMO_FetchXorTestType)
|
||||
? 1
|
||||
: 0;
|
||||
|
||||
if (fetch_op == 1) {
|
||||
ret = *std::max_element(_ret_val, _ret_val + args.num_wgs);
|
||||
} else {
|
||||
ret = *std::max_element(_s_buf, _s_buf + args.num_wgs);
|
||||
}
|
||||
if (ret != expected_val) {
|
||||
std::cerr << "data validation error\n";
|
||||
std::cerr << "got " << ret << ", expected " << expected_val << std::endl;
|
||||
exit(-1);
|
||||
}
|
||||
// PE 0 checks returns; target PE checks dest.
|
||||
if (args.myid) {
|
||||
verifyDestValues();
|
||||
} else {
|
||||
verifyReturnValues();
|
||||
}
|
||||
}
|
||||
|
||||
@@ -127,50 +247,51 @@ void AMOBitwiseTester<T>::verifyResults(size_t size) {
|
||||
template <> \
|
||||
__global__ void AMOBitwiseTest<T>( \
|
||||
int loop, int skip, long long int *start_time, \
|
||||
long long int *end_time, char *r_buf, T *s_buf, T *ret_val, \
|
||||
TestType type, ShmemContextType ctx_type) { \
|
||||
long long int *end_time, T *dest, T *ret_val, \
|
||||
AddrMode addr_mode, TestType type, ShmemContextType ctx_type) { \
|
||||
__shared__ rocshmem_ctx_t ctx; \
|
||||
int wg_id = get_flat_grid_id(); \
|
||||
int wg_id = get_flat_grid_id(); \
|
||||
int global_id = get_flat_id(); \
|
||||
int n_threads = get_flat_grid_size(); \
|
||||
int n_wgs = get_grid_num_blocks(); \
|
||||
rocshmem_wg_init(); \
|
||||
rocshmem_wg_ctx_create(ctx_type, &ctx); \
|
||||
if (hipThreadIdx_x == 0) { \
|
||||
for (int i = 0; i < loop + skip; i++) { \
|
||||
T *ptr = compute_target_ptr<T>(dest, addr_mode, wg_id, i, n_wgs); \
|
||||
T ret = 0; \
|
||||
T cond = 0; \
|
||||
for (int i = 0; i < loop + skip; i++) { \
|
||||
if (i == skip) { \
|
||||
start_time[wg_id] = wall_clock64(); \
|
||||
} \
|
||||
switch (type) { \
|
||||
case AMO_FetchAndTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_and(ctx, (T *)r_buf, \
|
||||
0xFFFF, 1); \
|
||||
break; \
|
||||
case AMO_AndTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_and(ctx, (T *)r_buf, 0xFFFF, 1); \
|
||||
break; \
|
||||
case AMO_FetchOrTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_or(ctx, (T *)r_buf, \
|
||||
0xFFFF, 1); \
|
||||
break; \
|
||||
case AMO_OrTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_or(ctx, (T *)r_buf, 0xFFFF, 1); \
|
||||
break; \
|
||||
case AMO_FetchXorTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_xor(ctx, (T *)r_buf, \
|
||||
0xFFFF, 1); \
|
||||
break; \
|
||||
case AMO_XorTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_xor(ctx, (T *)r_buf, 0xFFFF, 1); \
|
||||
break; \
|
||||
default: \
|
||||
break; \
|
||||
} \
|
||||
if (i == skip) { \
|
||||
start_time[wg_id] = wall_clock64(); \
|
||||
} \
|
||||
rocshmem_ctx_quiet(ctx); \
|
||||
end_time[wg_id] = wall_clock64(); \
|
||||
ret_val[wg_id] = ret; \
|
||||
rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \
|
||||
switch (type) { \
|
||||
case AMO_FetchAndTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_and(ctx, ptr, \
|
||||
(T)~(T)0, 1); \
|
||||
break; \
|
||||
case AMO_AndTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_and(ctx, ptr, (T)~(T)0, 1); \
|
||||
break; \
|
||||
case AMO_FetchOrTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_or(ctx, ptr, \
|
||||
(T)~(T)0, 1); \
|
||||
break; \
|
||||
case AMO_OrTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_or(ctx, ptr, (T)~(T)0, 1); \
|
||||
break; \
|
||||
case AMO_FetchXorTestType: \
|
||||
ret = rocshmem_ctx_##TNAME##_atomic_fetch_xor(ctx, ptr, \
|
||||
(T)~(T)0, 1); \
|
||||
break; \
|
||||
case AMO_XorTestType: \
|
||||
rocshmem_ctx_##TNAME##_atomic_xor(ctx, ptr, (T)~(T)0, 1); \
|
||||
break; \
|
||||
default: \
|
||||
break; \
|
||||
} \
|
||||
ret_val[global_id + i * n_threads] = ret; \
|
||||
} \
|
||||
rocshmem_ctx_quiet(ctx); \
|
||||
end_time[wg_id] = wall_clock64(); \
|
||||
__syncthreads(); \
|
||||
rocshmem_wg_ctx_destroy(&ctx); \
|
||||
rocshmem_wg_finalize(); \
|
||||
} \
|
||||
@@ -179,3 +300,5 @@ void AMOBitwiseTester<T>::verifyResults(size_t size) {
|
||||
AMO_BITWISE_DEF_GEN(unsigned int, uint)
|
||||
AMO_BITWISE_DEF_GEN(unsigned long, ulong)
|
||||
AMO_BITWISE_DEF_GEN(unsigned long long, ulonglong)
|
||||
AMO_BITWISE_DEF_GEN(int32_t, int32)
|
||||
AMO_BITWISE_DEF_GEN(int64_t, int64)
|
||||
|
||||
Reference in New Issue
Block a user