AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Merge pull request #44 from avinashkethineedi/fix/time_calculations

Sfoglia il codice sorgente 
Update bandwidth and latency calculations

[ROCm/rocshmem commit: f8701a44fa]
This commit is contained in:
Avinash Kethineedi
2025-02-17 12:48:33 -06:00
committato da GitHub
parent 04889cb71c 6c70aee32e
commit f2990df7f2
29 ha cambiato i file con 285 aggiunte e 289 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/rocshmem/include/rocshmem/rocshmem.hpp
-11
Vedi File
@@ -512,17 +512,6 @@ __device__ ATTR_NO_INLINE void rocshmem_wg_team_sync(rocshmem_team_t team);
*/
__device__ ATTR_NO_INLINE void *rocshmem_ptr(const void *dest, int pe);
/**
* @brief Query the current time. Similar to gettimeofday() on the CPU. To use
* this function, rocSHMEM must be configured with profiling support
* (--enable-profile).
*
* Can be called per thread with no performance penalty.
*
* @return Time in micro-seconds.
*/
__device__ uint64_t rocshmem_timer();
/**
* @brief Make all uncacheable GPU data visible to other agents in the sytem.
*
projects/rocshmem/src/rocshmem_gpu.cpp
-6
Vedi File
@@ -638,12 +638,6 @@ __device__ int rocshmem_my_pe() {
return get_internal_ctx(ROCSHMEM_CTX_DEFAULT)->my_pe;
}
__device__ uint64_t rocshmem_timer() {
GPU_DPRINTF("Function: rocshmem_timer\n");
return __read_clock();
}
template <typename T>
__device__ T rocshmem_atomic_fetch_add(rocshmem_ctx_t ctx, T *dest, T val,
int pe) {
projects/rocshmem/src/stats.hpp
+2 -2
Vedi File
@@ -141,10 +141,10 @@ class Stats {
StatType stats[I] = {0};
public:
__device__ uint64_t startTimer() const { return rocshmem_timer(); }
__device__ uint64_t startTimer() const { return wall_clock64(); }
__device__ void endTimer(uint64_t start, int index) {
incStat(index, rocshmem_timer() - start);
incStat(index, wall_clock64() - start);
}
__device__ void incStat(int index, int value = 1) {
projects/rocshmem/tests/functional_tests/alltoall_tester.cpp
+9 -8
Vedi File
@@ -58,10 +58,12 @@ rocshmem_team_t team_alltoall_world_dup;
* DEVICE TEST KERNEL
*****************************************************************************/
template <typename T1>
__global__ void AlltoallTest(int loop, int skip, uint64_t *timer,
T1 *source_buf, T1 *dest_buf, int size,
ShmemContextType ctx_type, rocshmem_team_t team) {
__global__ void AlltoallTest(int loop, int skip, long long int *start_time,
long long int *end_time, T1 *source_buf,
T1 *dest_buf, int size, ShmemContextType ctx_type,
rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -70,10 +72,9 @@ __global__ void AlltoallTest(int loop, int skip, uint64_t *timer,
__syncthreads();
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip && hipThreadIdx_x == 0) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
wg_alltoall<T1>(ctx, team,
dest_buf, // T* dest
@@ -84,7 +85,7 @@ __global__ void AlltoallTest(int loop, int skip, uint64_t *timer,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -126,8 +127,8 @@ void AlltoallTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(AlltoallTest<T1>, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, source_buf, dest_buf, size,
_shmem_context, team_alltoall_world_dup);
stream, loop, args.skip, start_time, end_time, source_buf,
dest_buf, size, _shmem_context, team_alltoall_world_dup);
num_msgs = loop + args.skip;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/amo_bitwise_tester.cpp
+13 -9
Vedi File
@@ -29,7 +29,8 @@ using namespace rocshmem;
/* Declare the global kernel template with a generic implementation */
template <typename T>
__global__ void AMOBitwiseTest(int loop, int skip, uint64_t *timer, char *r_buf,
__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,
ShmemContextType ctx_type) {
return;
@@ -64,8 +65,8 @@ 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, timer, _r_buf, _s_buf, _ret_val, _type,
_shmem_context);
loop, args.skip, start_time, end_time, _r_buf, _s_buf,
_ret_val, _type, _shmem_context);
_gridSize = gridsize;
num_msgs = (loop + args.skip) * gridsize.x;
@@ -123,17 +124,20 @@ void AMOBitwiseTester<T>::verifyResults(uint64_t size) {
#define AMO_BITWISE_DEF_GEN(T, TNAME) \
template <> \
__global__ void AMOBitwiseTest<T>( \
int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \
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) { \
__shared__ rocshmem_ctx_t ctx; \
int wg_id = get_flat_grid_id(); \
rocshmem_wg_init(); \
rocshmem_wg_ctx_create(ctx_type, &ctx); \
if (hipThreadIdx_x == 0) { \
uint64_t start; \
T ret = 0; \
T cond = 0; \
for (int i = 0; i < loop + skip; i++) { \
if (i == skip) start = rocshmem_timer(); \
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, \
@@ -161,9 +165,9 @@ void AMOBitwiseTester<T>::verifyResults(uint64_t size) {
} \
} \
rocshmem_ctx_quiet(ctx); \
timer[hipBlockIdx_x] = rocshmem_timer() - start; \
ret_val[hipBlockIdx_x] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \
end_time[wg_id] = wall_clock64(); \
ret_val[wg_id] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \
} \
rocshmem_wg_ctx_destroy(&ctx); \
rocshmem_wg_finalize(); \
projects/rocshmem/tests/functional_tests/amo_extended_tester.cpp
+15 -11
Vedi File
@@ -29,9 +29,10 @@ using namespace rocshmem;
/* Declare the global kernel template with a generic implementation */
template <typename T>
__global__ void AMOExtendedTest(int loop, int skip, uint64_t *timer,
char *r_buf, T *s_buf, T *ret_val,
TestType type, ShmemContextType ctx_type) {
__global__ void AMOExtendedTest(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) {
return;
}
@@ -64,8 +65,8 @@ void AMOExtendedTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(AMOExtendedTest, gridsize, blocksize, shared_bytes, stream,
loop, args.skip, timer, _r_buf, _s_buf, _ret_val, _type,
_shmem_context);
loop, args.skip, start_time, end_time, _r_buf, _s_buf,
_ret_val, _type, _shmem_context);
_gridSize = gridsize;
num_msgs = (loop + args.skip) * gridsize.x;
@@ -111,17 +112,20 @@ void AMOExtendedTester<T>::verifyResults(uint64_t size) {
#define AMO_EXTENDED_DEF_GEN(T, TNAME) \
template <> \
__global__ void AMOExtendedTest<T>( \
int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \
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) { \
__shared__ rocshmem_ctx_t ctx; \
int wg_id = get_flat_grid_id(); \
rocshmem_wg_init(); \
rocshmem_wg_ctx_create(ctx_type, &ctx); \
if (hipThreadIdx_x == 0) { \
uint64_t start; \
T ret = 0; \
T cond = 0; \
for (int i = 0; i < loop + skip; i++) { \
if (i == skip) start = rocshmem_timer(); \
if (i == skip) { \
start_time[wg_id] = wall_clock64(); \
} \
switch (type) { \
case AMO_FetchTestType: \
ret = rocshmem_ctx_##TNAME##_atomic_fetch(ctx, (T *)r_buf, 1); \
@@ -138,9 +142,9 @@ void AMOExtendedTester<T>::verifyResults(uint64_t size) {
} \
} \
rocshmem_ctx_quiet(ctx); \
timer[hipBlockIdx_x] = rocshmem_timer() - start; \
ret_val[hipBlockIdx_x] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \
end_time[wg_id] = wall_clock64(); \
ret_val[wg_id] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \
} \
rocshmem_wg_ctx_destroy(&ctx); \
rocshmem_wg_finalize(); \
projects/rocshmem/tests/functional_tests/amo_standard_tester.cpp
+15 -11
Vedi File
@@ -29,9 +29,10 @@ using namespace rocshmem;
/* Declare the global kernel template with a generic implementation */
template <typename T>
__global__ void AMOStandardTest(int loop, int skip, uint64_t *timer,
char *r_buf, T *s_buf, T *ret_val,
TestType type, ShmemContextType ctx_type) {
__global__ void AMOStandardTest(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) {
return;
}
@@ -64,8 +65,8 @@ void AMOStandardTester<T>::launchKernel(dim3 gridsize, dim3 blocksize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(AMOStandardTest, gridsize, blocksize, shared_bytes, stream,
loop, args.skip, timer, _r_buf, _s_buf, _ret_val, _type,
_shmem_context);
loop, args.skip, start_time, end_time, _r_buf, _s_buf,
_ret_val, _type, _shmem_context);
_gridSize = gridsize;
num_msgs = (loop + args.skip) * gridsize.x;
@@ -119,17 +120,20 @@ void AMOStandardTester<T>::verifyResults(uint64_t size) {
#define AMO_STANDARD_DEF_GEN(T, TNAME) \
template <> \
__global__ void AMOStandardTest<T>( \
int loop, int skip, uint64_t *timer, char *r_buf, T *s_buf, T *ret_val, \
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) { \
__shared__ rocshmem_ctx_t ctx; \
int wg_id = get_flat_grid_id(); \
rocshmem_wg_init(); \
rocshmem_wg_ctx_create(ctx_type, &ctx); \
if (hipThreadIdx_x == 0) { \
uint64_t start; \
T ret = 0; \
T cond = 0; \
for (int i = 0; i < loop + skip; i++) { \
if (i == skip) start = rocshmem_timer(); \
if (i == skip) { \
start_time[wg_id] = wall_clock64(); \
} \
switch (type) { \
case AMO_FAddTestType: \
ret = rocshmem_ctx_##TNAME##_atomic_fetch_add(ctx, (T *)r_buf, 2, \
@@ -155,9 +159,9 @@ void AMOStandardTester<T>::verifyResults(uint64_t size) {
} \
} \
rocshmem_ctx_quiet(ctx); \
timer[hipBlockIdx_x] = rocshmem_timer() - start; \
ret_val[hipBlockIdx_x] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[hipBlockIdx_x], r_buf, sizeof(T), 1); \
end_time[wg_id] = wall_clock64(); \
ret_val[wg_id] = ret; \
rocshmem_ctx_getmem(ctx, &s_buf[wg_id], r_buf, sizeof(T), 1); \
} \
rocshmem_wg_ctx_destroy(&ctx); \
rocshmem_wg_finalize(); \
projects/rocshmem/tests/functional_tests/barrier_all_tester.cpp
+6 -5
Vedi File
@@ -29,16 +29,17 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void BarrierAllTest(int loop, int skip, uint64_t *timer) {
__global__ void BarrierAllTest(int loop, int skip, long long int *start_time,
long long int *end_time) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ROCSHMEM_CTX_WG_PRIVATE, &ctx);
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (hipThreadIdx_x == 0 && i == skip) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
__syncthreads();
@@ -48,7 +49,7 @@ __global__ void BarrierAllTest(int loop, int skip, uint64_t *timer) {
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -67,7 +68,7 @@ void BarrierAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(BarrierAllTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer);
loop, args.skip, start_time, end_time);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/empty_tester.cpp
+6 -4
Vedi File
@@ -29,8 +29,9 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void EmptyTest(int loop, int skip, uint64_t *timer, int size,
TestType type, ShmemContextType ctx_type) {
__global__ void EmptyTest(int loop, int skip, long long int *start_time,
long long int *end_time, int size, TestType type,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -52,8 +53,9 @@ void EmptyTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
uint64_t size) {
size_t shared_bytes = 0;
hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream, loop,
args.skip, timer, size, _type, _shmem_context);
hipLaunchKernelGGL(EmptyTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, start_time, end_time, size, _type,
_shmem_context);
}
void EmptyTester::verifyResults(uint64_t size) {}
projects/rocshmem/tests/functional_tests/extended_primitives.cpp
+11 -8
Vedi File
@@ -31,11 +31,13 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
char *s_buf, char *r_buf, int size,
TestType type,
__global__ void ExtendedPrimitiveTest(int loop, int skip,
long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, TestType type,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -44,13 +46,14 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
* If the number of work groups is greater than 1, this kernel performs a
* tiled functional test
*/
uint64_t start;
uint64_t idx = size * get_flat_grid_id();
s_buf += idx;
r_buf += idx;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) start = rocshmem_timer();
if (i == skip) {
start_time[wg_id] = wall_clock64();
}
switch (type) {
case WGGetTestType:
@@ -73,7 +76,7 @@ __global__ void ExtendedPrimitiveTest(int loop, int skip, uint64_t *timer,
rocshmem_ctx_quiet(ctx);
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -105,8 +108,8 @@ void ExtendedPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(ExtendedPrimitiveTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, (char*)s_buf,
(char*)r_buf, size, _type, _shmem_context);
stream, loop, args.skip, start_time, end_time,
(char*)s_buf, (char*)r_buf, size, _type, _shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop * gridSize.x;
projects/rocshmem/tests/functional_tests/fcollect_tester.cpp
+9 -8
Vedi File
@@ -58,10 +58,12 @@ FCOLLECT_DEF_GEN(unsigned long long, ulonglong)
* DEVICE TEST KERNEL
*****************************************************************************/
template <typename T1>
__global__ void FcollectTest(int loop, int skip, uint64_t *timer,
T1 *source_buf, T1 *dest_buf, int size,
ShmemContextType ctx_type, rocshmem_team_t team) {
__global__ void FcollectTest(int loop, int skip, long long int *start_time,
long long int *end_time, T1 *source_buf,
T1 *dest_buf, int size, ShmemContextType ctx_type,
rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -69,10 +71,9 @@ __global__ void FcollectTest(int loop, int skip, uint64_t *timer,
int n_pes = rocshmem_ctx_n_pes(ctx);
__syncthreads();
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip && hipThreadIdx_x == 0) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
wg_fcollect<T1>(ctx, team,
dest_buf, // T* dest
@@ -83,7 +84,7 @@ __global__ void FcollectTest(int loop, int skip, uint64_t *timer,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -125,8 +126,8 @@ void FcollectTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(FcollectTest<T1>, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, source_buf, dest_buf, size,
_shmem_context, team_fcollect_world_dup);
stream, loop, args.skip, start_time, end_time, source_buf,
dest_buf, size, _shmem_context, team_fcollect_world_dup);
num_msgs = loop + args.skip;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/ping_all_tester.cpp
+7 -5
Vedi File
@@ -29,9 +29,11 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf,
__global__ void PingAllTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *r_buf,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -44,12 +46,11 @@ __global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf,
}
if (hipThreadIdx_x == 0) {
uint64_t start;
auto blk_pe_off {hipBlockIdx_x * num_pe};
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
for (int j{0}; j < num_pe; j++) {
@@ -57,7 +58,7 @@ __global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf,
}
rocshmem_int_wait_until_all(&r_buf[blk_pe_off], num_pe, status, ROCSHMEM_CMP_EQ, 1);
}
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
rocshmem_wg_finalize();
@@ -83,7 +84,8 @@ void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(PingAllTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, r_buf, _shmem_context);
loop, args.skip, start_time, end_time, r_buf,
_shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/ping_all_tester.hpp
+2 -1
Vedi File
@@ -28,7 +28,8 @@
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PingAllTest(int loop, int skip, uint64_t *timer, int *r_buf);
__global__ void PingAllTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *r_buf);
/******************************************************************************
* HOST TESTER CLASS
projects/rocshmem/tests/functional_tests/ping_pong_tester.cpp
+7 -5
Vedi File
@@ -29,9 +29,11 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf,
__global__ void PingPongTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *r_buf,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -39,11 +41,10 @@ __global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf,
int pe = rocshmem_ctx_my_pe(ctx);
if (hipThreadIdx_x == 0) {
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
if (pe == 0) {
@@ -56,7 +57,7 @@ __global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf,
rocshmem_ctx_int_p(ctx, &r_buf[hipBlockIdx_x], i + 1, 0);
}
}
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
rocshmem_wg_finalize();
@@ -80,7 +81,8 @@ void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(PingPongTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, r_buf, _shmem_context);
loop, args.skip, start_time, end_time, r_buf,
_shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/ping_pong_tester.hpp
+2 -1
Vedi File
@@ -28,7 +28,8 @@
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PingPongTest(int loop, int skip, uint64_t *timer, int *r_buf);
__global__ void PingPongTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *r_buf);
/******************************************************************************
* HOST TESTER CLASS
projects/rocshmem/tests/functional_tests/primitive_mr_tester.cpp
+9 -6
Vedi File
@@ -29,17 +29,18 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PrimitiveMRTest(int loop, uint64_t *timer, char *s_buf,
__global__ void PrimitiveMRTest(int loop, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
if (hipThreadIdx_x == 0) {
uint64_t start;
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
for (int win_i = 0; win_i < 64 * loop; win_i++) {
for (int i = 0; i < 64; i++) {
@@ -48,7 +49,7 @@ __global__ void PrimitiveMRTest(int loop, uint64_t *timer, char *s_buf,
rocshmem_ctx_quiet(ctx);
}
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
__syncthreads();
@@ -81,11 +82,13 @@ void PrimitiveMRTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
/* Warmup */
hipLaunchKernelGGL(PrimitiveMRTest, gridSize, blockSize, shared_bytes, stream,
loop, timer, s_buf, r_buf, size, _shmem_context);
loop, start_time, end_time, s_buf, r_buf, size,
_shmem_context);
/* Benchmark */
hipLaunchKernelGGL(PrimitiveMRTest, gridSize, blockSize, shared_bytes, stream,
loop, timer, s_buf, r_buf, size, _shmem_context);
loop, start_time, end_time, s_buf, r_buf, size,
_shmem_context);
CHECK_HIP(hipDeviceSynchronize());
projects/rocshmem/tests/functional_tests/primitive_tester.cpp
+7 -7
Vedi File
@@ -29,19 +29,19 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PrimitiveTest(int loop, int skip, uint64_t *timer, char *s_buf,
__global__ void PrimitiveTest(int loop, int skip, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, TestType type,
ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
__syncthreads();
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
switch (type) {
@@ -79,7 +79,7 @@ __global__ void PrimitiveTest(int loop, int skip, uint64_t *timer, char *s_buf,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -109,8 +109,8 @@ void PrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(PrimitiveTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, s_buf, r_buf, size, _type,
_shmem_context);
loop, args.skip, start_time, end_time, s_buf, r_buf,
size, _type, _shmem_context);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/random_access_tester.cpp
+14 -9
Vedi File
@@ -49,13 +49,14 @@ __device__ bool thread_passing(int num_bins, uint32_t *bin_threads,
return pass;
}
__global__ void RandomAccessTest(int loop, int skip, uint64_t *timer,
int *s_buf, int *r_buf, int size, OpType type,
__global__ void RandomAccessTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *s_buf,
int *r_buf, int size, OpType type,
int coal_coef, int num_bins, int num_waves,
uint32_t *threads_bins, uint32_t *off_bins,
uint32_t *PE_bins, ShmemContextType ctx_type) {
uint64_t start;
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -69,7 +70,9 @@ __global__ void RandomAccessTest(int loop, int skip, uint64_t *timer,
r_buf = r_buf + offset;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) start = rocshmem_timer();
if (i == skip) {
start_time[wg_id] = wall_clock64();
}
switch (type) {
case GetType:
rocshmem_ctx_getmem(ctx, r_buf, s_buf, size, PE);
@@ -84,8 +87,9 @@ __global__ void RandomAccessTest(int loop, int skip, uint64_t *timer,
rocshmem_ctx_quiet(ctx);
atomicAdd((unsigned long long *)&timer[hipBlockIdx_x],
rocshmem_timer() - start);
// atomicAdd((unsigned long long *)&timer[hipBlockIdx_x],
// rocshmem_timer() - start);
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
rocshmem_wg_finalize();
@@ -178,9 +182,10 @@ void RandomAccessTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
if (args.myid == 0) {
hipLaunchKernelGGL(RandomAccessTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, s_buf, r_buf, size,
(OpType)args.op_type, _coal_coef, _num_bins, _num_waves,
_threads_bins, _off_bins, _PE_bins, _shmem_context);
stream, loop, args.skip, start_time, end_time, s_buf,
r_buf, size, (OpType)args.op_type, _coal_coef,
_num_bins, _num_waves, _threads_bins, _off_bins,
_PE_bins, _shmem_context);
}
num_msgs = (loop + args.skip) * _num_waves * _thread_access;
num_timed_msgs = loop * _num_waves * _thread_access;
projects/rocshmem/tests/functional_tests/random_access_tester.hpp
+3 -2
Vedi File
@@ -28,8 +28,9 @@
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void RandomAccessTest(int loop, int skip, uint64_t *timer,
int *s_buf, int *r_buf, int size, OpType type,
__global__ void RandomAccessTest(int loop, int skip, long long int *start_time,
long long int *end_time, int *s_buf,
int *r_buf, int size, OpType type,
int coal_coef, int num_bins, int num_waves,
uint32_t *threads_bins, uint32_t *off_bins,
uint32_t *PE_bins);
projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp
+12 -10
Vedi File
@@ -29,20 +29,21 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void PutmemSignalTest(int loop, int skip, uint64_t *timer, char *s_buf,
__global__ void PutmemSignalTest(int loop, int skip, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, uint64_t *sig_addr,
TestType type, ShmemContextType ctx_type, int sig_op) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
uint64_t start;
uint64_t signal = 1;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
__syncthreads();
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
switch (type) {
@@ -74,23 +75,24 @@ __global__ void PutmemSignalTest(int loop, int skip, uint64_t *timer, char *s_bu
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
rocshmem_wg_finalize();
}
__global__ void SignalFetchTest(int loop, int skip, uint64_t *timer, uint64_t *sig_addr,
__global__ void SignalFetchTest(int loop, int skip, long long int *start_time,
long long int *end_time, uint64_t *sig_addr,
uint64_t *fetched_value, TestType type) {
rocshmem_wg_init();
uint64_t start;
int wg_id = get_flat_grid_id();
for (int i = 0; i < loop + skip; i++) {
if (i == skip) {
__syncthreads();
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
switch (type) {
@@ -111,7 +113,7 @@ __global__ void SignalFetchTest(int loop, int skip, uint64_t *timer, uint64_t *s
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_finalize();
@@ -155,10 +157,10 @@ void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int
(_type == WAVESignalFetchTestType) ||
(_type == WGSignalFetchTestType)) {
hipLaunchKernelGGL(SignalFetchTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, sig_addr, fetched_value, _type);
loop, args.skip, start_time, end_time, sig_addr, fetched_value, _type);
} else {
hipLaunchKernelGGL(PutmemSignalTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, s_buf, r_buf, size, sig_addr,
loop, args.skip, start_time, end_time, s_buf, r_buf, size, sig_addr,
_type, _shmem_context, sig_op);
}
projects/rocshmem/tests/functional_tests/swarm_tester.cpp
+12 -8
Vedi File
@@ -29,9 +29,11 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf,
__global__ void GetSwarmTest(int loop, int skip, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, ShmemContextType ctx_type) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
int provided;
rocshmem_wg_init_thread(ROCSHMEM_THREAD_MULTIPLE, &provided);
@@ -42,18 +44,20 @@ __global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf,
__syncthreads();
int index = hipThreadIdx_x * size;
uint64_t start = 0;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) start = rocshmem_timer();
if (i == skip) {
start_time[wg_id] = wall_clock64();
}
rocshmem_ctx_getmem(ctx, &r_buf[index], &s_buf[index], size, 1);
__syncthreads();
}
atomicAdd((unsigned long long *)&timer[hipBlockIdx_x],
rocshmem_timer() - start);
// atomicAdd((unsigned long long *)&timer[hipBlockIdx_x],
// rocshmem_timer() - start);
end_time[wg_id] = wall_clock64();
rocshmem_wg_ctx_destroy(&ctx);
rocshmem_wg_finalize();
@@ -71,8 +75,8 @@ void GetSwarmTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
size_t shared_bytes = 0;
hipLaunchKernelGGL(GetSwarmTest, gridSize, blockSize, shared_bytes, stream,
loop, args.skip, timer, s_buf, r_buf, size,
_shmem_context);
loop, args.skip, start_time, end_time, s_buf, r_buf, size,
_shmem_context);
num_msgs = (loop + args.skip) * gridSize.x * blockSize.x;
num_timed_msgs = loop * gridSize.x * blockSize.x;
projects/rocshmem/tests/functional_tests/swarm_tester.hpp
+2 -1
Vedi File
@@ -28,7 +28,8 @@
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void GetSwarmTest(int loop, int skip, uint64_t *timer, char *s_buf,
__global__ void GetSwarmTest(int loop, int skip, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size);
/******************************************************************************
projects/rocshmem/tests/functional_tests/sync_tester.cpp
+7 -5
Vedi File
@@ -30,16 +30,18 @@ rocshmem_team_t team_sync_world_dup;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void SyncTest(int loop, int skip, uint64_t *timer, TestType type,
__global__ void SyncTest(int loop, int skip, long long int *start_time,
long long int *end_time, TestType type,
ShmemContextType ctx_type, rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (hipThreadIdx_x == 0 && i == skip) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
__syncthreads();
@@ -57,7 +59,7 @@ __global__ void SyncTest(int loop, int skip, uint64_t *timer, TestType type,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -84,7 +86,7 @@ void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,
&team_sync_world_dup);
hipLaunchKernelGGL(SyncTest, gridSize, blockSize, shared_bytes, stream, loop,
args.skip, timer, _type, _shmem_context,
args.skip, start_time, end_time, _type, _shmem_context,
team_sync_world_dup);
num_msgs = (loop + args.skip) * gridSize.x;
projects/rocshmem/tests/functional_tests/team_broadcast_tester.cpp
+8 -7
Vedi File
@@ -61,11 +61,13 @@ rocshmem_team_t team_bcast_world_dup;
* DEVICE TEST KERNEL
*****************************************************************************/
template <typename T1>
__global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer,
T1 *source_buf, T1 *dest_buf, int size,
__global__ void TeamBroadcastTest(int loop, int skip, long long int *start_time,
long long int *end_time, T1 *source_buf,
T1 *dest_buf, int size,
ShmemContextType ctx_type,
rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -74,10 +76,9 @@ __global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer,
__syncthreads();
uint64_t start;
for (int i = 0; i < loop; i++) {
if (i == skip && hipThreadIdx_x == 0) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
wg_team_broadcast<T1>(ctx, team,
@@ -91,7 +92,7 @@ __global__ void TeamBroadcastTest(int loop, int skip, uint64_t *timer,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -131,8 +132,8 @@ void TeamBroadcastTester<T1>::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(TeamBroadcastTest<T1>, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, source_buf, dest_buf, size,
_shmem_context, team_bcast_world_dup);
stream, loop, args.skip, start_time, end_time, source_buf,
dest_buf, size, _shmem_context, team_bcast_world_dup);
num_msgs = loop + args.skip;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/team_ctx_primitive_tester.cpp
+13 -9
Vedi File
@@ -31,20 +31,23 @@ rocshmem_team_t team_primitive_world_dup;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void TeamCtxPrimitiveTest(int loop, int skip, uint64_t *timer,
char *s_buf, char *r_buf, int size,
TestType type, ShmemContextType ctx_type,
__global__ void TeamCtxPrimitiveTest(int loop, int skip, long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, TestType type,
ShmemContextType ctx_type,
rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_team_create_ctx(team, ctx_type, &ctx);
if (hipThreadIdx_x == 0) {
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) start = rocshmem_timer();
if (i == skip) {
start_time[wg_id] = wall_clock64();
}
switch (type) {
case TeamCtxGetTestType:
rocshmem_ctx_getmem(ctx, r_buf, s_buf, size, 1);
@@ -65,7 +68,7 @@ __global__ void TeamCtxPrimitiveTest(int loop, int skip, uint64_t *timer,
rocshmem_ctx_quiet(ctx);
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -104,8 +107,9 @@ void TeamCtxPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(TeamCtxPrimitiveTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, s_buf, r_buf, size, _type,
_shmem_context, team_primitive_world_dup);
stream, loop, args.skip, start_time, end_time, s_buf,
r_buf, size, _type, _shmem_context,
team_primitive_world_dup);
num_msgs = (loop + args.skip) * gridSize.x;
num_timed_msgs = loop * gridSize.x;
projects/rocshmem/tests/functional_tests/team_reduction_tester.cpp
+9 -8
Vedi File
@@ -73,11 +73,13 @@ rocshmem_team_t team_reduce_world_dup;
* DEVICE TEST KERNEL
*****************************************************************************/
template <typename T1, ROCSHMEM_OP T2>
__global__ void TeamReductionTest(int loop, int skip, uint64_t *timer,
T1 *s_buf, T1 *r_buf, int size, TestType type,
__global__ void TeamReductionTest(int loop, int skip, long long int *start_time,
long long int *end_time, T1 *s_buf, T1 *r_buf,
int size, TestType type,
ShmemContextType ctx_type,
rocshmem_team_t team) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -86,10 +88,9 @@ __global__ void TeamReductionTest(int loop, int skip, uint64_t *timer,
__syncthreads();
uint64_t start;
for (int i = 0; i < loop + skip; i++) {
if (i == skip && hipThreadIdx_x == 0) {
start = rocshmem_timer();
start_time[wg_id] = wall_clock64();
}
wg_team_reduce<T1, T2>(ctx, team, r_buf, s_buf, size);
rocshmem_ctx_wg_barrier_all(ctx);
@@ -98,7 +99,7 @@ __global__ void TeamReductionTest(int loop, int skip, uint64_t *timer,
__syncthreads();
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[wg_id] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -138,9 +139,9 @@ void TeamReductionTester<T1, T2>::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(HIP_KERNEL_NAME(TeamReductionTest<T1, T2>), gridSize,
blockSize, shared_bytes, stream, loop, args.skip, timer,
s_buf, r_buf, size, _type, _shmem_context,
team_reduce_world_dup);
blockSize, shared_bytes, stream, loop, args.skip,
start_time, end_time, s_buf, r_buf, size, _type,
_shmem_context, team_reduce_world_dup);
num_msgs = loop + args.skip;
num_timed_msgs = loop;
projects/rocshmem/tests/functional_tests/tester.cpp
+59 -108
Vedi File
@@ -58,14 +58,22 @@ Tester::Tester(TesterArguments args) : args(args) {
_shmem_context = args.shmem_context;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
num_warps = (args.wg_size - 1) / deviceProps.warpSize + 1;
wf_size = deviceProps.warpSize;
num_warps = (args.wg_size - 1) / wf_size + 1;
CHECK_HIP(hipStreamCreate(&stream));
CHECK_HIP(hipEventCreate(&start_event));
CHECK_HIP(hipEventCreate(&stop_event));
CHECK_HIP(hipMalloc((void**)&timer, sizeof(uint64_t) * args.num_wgs));
CHECK_HIP(hipDeviceGetAttribute(&wall_clk_rate,
hipDeviceAttributeWallClockRate, device_id));
num_timers = args.num_wgs;
CHECK_HIP(hipMalloc((void**)&timer, sizeof(long long int) * num_timers));
CHECK_HIP(hipMalloc((void**)&start_time, sizeof(long long int) * num_timers));
CHECK_HIP(hipMalloc((void**)&end_time, sizeof(long long int) * num_timers));
}
Tester::~Tester() {
CHECK_HIP(hipFree(end_time));
CHECK_HIP(hipFree(start_time));
CHECK_HIP(hipFree(timer));
CHECK_HIP(hipEventDestroy(stop_event));
CHECK_HIP(hipEventDestroy(start_event));
@@ -75,11 +83,6 @@ Tester::~Tester() {
std::vector<Tester*> Tester::create(TesterArguments args) {
int rank = args.myid;
std::vector<Tester*> testers;
hipDeviceProp_t deviceProps;
int device_id, numWarps;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
numWarps = (args.wg_size - 1) / deviceProps.warpSize + 1;
if (rank == 0) std::cout << "### Creating Test: ";
@@ -352,35 +355,23 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
testers.push_back(new ShmemPtrTester(args));
return testers;
case WGGetTestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Blocking WG level Gets ###" << std::endl;
else std::cout << "Blocking WG level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WG level Gets ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGGetNBITestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Non-Blocking WG level Gets ###" << std::endl;
else std::cout << "Non-Blocking WG level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WG level Gets ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGPutTestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Blocking WG level Puts ###" << std::endl;
else std::cout << "Blocking WG level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WG level Puts ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGPutNBITestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Non-Blocking WG level Puts ###" << std::endl;
else std::cout << "Non-Blocking WG level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WG level Puts ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case PutNBIMRTestType:
@@ -389,35 +380,23 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
testers.push_back(new PrimitiveMRTester(args));
return testers;
case WAVEGetTestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Gets ###" << std::endl;
else std::cout << "Blocking WAVE level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WAVE level Gets ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEGetNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Gets ###" << std::endl;
else std::cout << "Non-Blocking WAVE level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WAVE level Gets ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEPutTestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Puts ###" << std::endl;
else std::cout << "Blocking WAVE level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WAVE level Puts ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEPutNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Puts ###" << std::endl;
else std::cout << "Non-Blocking WAVE level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WAVE level Puts ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case PutSignalTestType:
@@ -501,11 +480,6 @@ void Tester::execute() {
* rocshmem pes.
*/
if (peLaunchesKernel()) {
/**
* TODO:
* Verify that this timer type is actually uint64_t on the
* device side.
*/
memset(timer, 0, sizeof(uint64_t) * args.num_wgs);
const dim3 blockSize(args.wg_size, 1, 1);
@@ -521,9 +495,6 @@ void Tester::execute() {
if (err != hipSuccess) {
printf("error = %d \n", err);
}
// rocshmem_dump_stats();
// rocshmem_reset_stats();
}
barrier();
@@ -533,32 +504,10 @@ void Tester::execute() {
// data validation
verifyResults(size);
/**
* Adjust size for *_wg and *_wave functions
*/
uint64_t size_ = size;
TestType type = (TestType)args.algorithm;
switch (type) {
case WAVEGetTestType:
case WAVEGetNBITestType:
case WAVEPutTestType:
case WAVEPutNBITestType:
size_ *= (args.num_wgs * num_warps);
break;
case WGGetTestType:
case WGGetNBITestType:
case WGPutTestType:
case WGPutNBITestType:
size_ *= args.num_wgs;
break;
default:
break;
}
barrier();
if (_type != TeamCtxInfraTestType) {
print(size_);
print(size);
}
}
}
@@ -589,30 +538,39 @@ void Tester::print(uint64_t size) {
return;
}
uint64_t timer_avg = timerAvgInMicroseconds();
double latency_avg = static_cast<double>(timer_avg) / num_timed_msgs;
/**
* Calculate total amount of data transfered
*/
uint64_t total_size = size * num_timed_msgs;
double timer_avg = timerAvgInMicroseconds();
double latency_avg = timer_avg / num_timed_msgs;
double avg_msg_rate = num_timed_msgs / (timer_avg / 1e6);
float total_kern_time_ms;
CHECK_HIP(hipEventElapsedTime(&total_kern_time_ms, start_event, stop_event));
float total_kern_time_s = total_kern_time_ms / 1000;
double time_us = gpuCyclesToMicroseconds(max_end_time - min_start_time);
double time_s = time_us / 1e6;
double bandwidth_avg_gbs =
num_msgs * size * bw_factor / total_kern_time_s / pow(2, 30);
static_cast<double>(total_size * bw_factor) / time_s / pow(2, 30);
int field_width = 20;
int float_precision = 2;
if (_print_header) {
printf("%-*s%*s%*s%*s",
10, "# Size (B)",
printf("%-*s%-*s%*s%*s%*s",
15, "# Size (B)",
15, "# of timed Msgs",
field_width, "Latency (us)",
field_width, "Bandwidth (GB/s)",
field_width + 1, "Msg Rate (Msg/s)\n");
_print_header = 0;
}
printf("%-*lu%*.*f%*.*f%*.*f\n",
10, size,
printf("%-*lu%-*d%*.*f%*.*f%*.*f\n",
15, size,
15, num_timed_msgs,
field_width, float_precision, latency_avg,
field_width, float_precision, bandwidth_avg_gbs,
field_width, float_precision, avg_msg_rate);
@@ -634,33 +592,26 @@ void Tester::barrier() {
flush_hdp();
}
uint64_t Tester::gpuCyclesToMicroseconds(uint64_t cycles) {
/**
* The dGPU asm core timer runs at 27MHz. This is different from the
* core clock returned by HIP. For an APU, this is different and might
* need adjusting.
*/
uint64_t gpu_frequency_MHz = 27;
/**
* hipDeviceGetAttribute(&gpu_frequency_khz,
* hipDeviceAttributeClockRate,
* 0);
*/
return cycles / gpu_frequency_MHz;
double Tester::gpuCyclesToMicroseconds(long long int cycles) {
return static_cast<double>(cycles) /
(static_cast<double>(wall_clk_rate) * 1e-3);
}
uint64_t Tester::timerAvgInMicroseconds() {
uint64_t sum = 0;
double Tester::timerAvgInMicroseconds() {
double sum = 0;
min_start_time = LLONG_MAX;
max_end_time = 0;
/**
* TODO: (bpotter/avinash) Modify the calcuation for the Tiled version of
* puts and gets at wavefront level
*/
for (uint64_t i = 0; i < args.num_wgs; i++) {
for (uint32_t i = 0; i < num_timers; i++) {
timer[i] = end_time[i] - start_time[i];
sum += gpuCyclesToMicroseconds(timer[i]);
min_start_time = (start_time[i] < min_start_time)
? start_time[i]
: min_start_time;
max_end_time = (end_time[i] > max_end_time)
? end_time[i]
: max_end_time;
}
return sum / args.num_wgs;
return sum / num_timers;
}
projects/rocshmem/tests/functional_tests/tester.hpp
+12 -3
Vedi File
@@ -25,8 +25,10 @@
#include <rocshmem/rocshmem.hpp>
#include <vector>
#include <climits>
#include "tester_arguments.hpp"
#include "../src/util.hpp"
/******************************************************************************
* TESTER CLASS TYPES
@@ -126,6 +128,8 @@ class Tester {
int num_warps = 0;
int bw_factor = 1;
int device_id = 0;
int wall_clk_rate = 0; //in kilohertz
int wf_size = 0;
TesterArguments args;
@@ -135,7 +139,12 @@ class Tester {
hipStream_t stream;
hipDeviceProp_t deviceProps;
uint64_t *timer = nullptr;
long long int *timer = nullptr;
long long int *start_time = nullptr;
long long int *end_time = nullptr;
long long int min_start_time = 0;
long long int max_end_time = 0;
uint32_t num_timers = 0;
private:
bool _print_header = 1;
@@ -143,9 +152,9 @@ class Tester {
void barrier();
uint64_t gpuCyclesToMicroseconds(uint64_t cycles);
double gpuCyclesToMicroseconds(long long int cycles);
uint64_t timerAvgInMicroseconds();
double timerAvgInMicroseconds();
bool peLaunchesKernel();
projects/rocshmem/tests/functional_tests/wave_level_primitives.cpp
+14 -11
Vedi File
@@ -31,11 +31,14 @@ using namespace rocshmem;
/******************************************************************************
* DEVICE TEST KERNEL
*****************************************************************************/
__global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
char *s_buf, char *r_buf, int size,
TestType type, ShmemContextType ctx_type,
int wf_size) {
__global__ void WaveLevelPrimitiveTest(int loop, int skip,
long long int *start_time,
long long int *end_time, char *s_buf,
char *r_buf, int size, TestType type,
ShmemContextType ctx_type, int wf_size) {
__shared__ rocshmem_ctx_t ctx;
int wg_id = get_flat_grid_id();
rocshmem_wg_init();
rocshmem_wg_ctx_create(ctx_type, &ctx);
@@ -44,7 +47,6 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
* If the number of wavefronts is greater than 1, this kernel performs a
* tiled functional test
*/
uint64_t start;
int wf_id = get_flat_block_id() / wf_size;
int wg_offset = size * get_flat_grid_id() * (get_flat_block_size() / wf_size);
int idx = wf_id * size + wg_offset;
@@ -52,8 +54,9 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
r_buf += idx;
for (int i = 0; i < loop + skip; i++) {
if (i == skip) start = rocshmem_timer();
if (i == skip) {
start_time[wg_id] = wall_clock64();
}
switch (type) {
case WAVEGetTestType:
rocshmem_ctx_getmem_wave(ctx, r_buf, s_buf, size, 1);
@@ -75,7 +78,7 @@ __global__ void WaveLevelPrimitiveTest(int loop, int skip, uint64_t *timer,
rocshmem_ctx_quiet(ctx);
if (hipThreadIdx_x == 0) {
timer[hipBlockIdx_x] = rocshmem_timer() - start;
end_time[hipBlockIdx_x] = wall_clock64();
}
rocshmem_wg_ctx_destroy(&ctx);
@@ -109,9 +112,9 @@ void WaveLevelPrimitiveTester::launchKernel(dim3 gridSize, dim3 blockSize,
size_t shared_bytes = 0;
hipLaunchKernelGGL(WaveLevelPrimitiveTest, gridSize, blockSize, shared_bytes,
stream, loop, args.skip, timer, (char*)s_buf,
(char*)r_buf, size, _type, _shmem_context,
deviceProps.warpSize);
stream, loop, args.skip, start_time, end_time,
(char*)s_buf, (char*)r_buf, size, _type, _shmem_context,
wf_size);
num_msgs = (loop + args.skip) * gridSize.x * num_warps;
num_timed_msgs = loop * gridSize.x * num_warps;