Konstantin Pyzhov
149 satır
4.0 KiB
C++
149 satır
4.0 KiB
C++
/* HIT_START
|
|
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
|
|
* RUN_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
|
|
* RUN_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
|
|
* RUN_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
|
|
* HIT_END
|
|
*/
|
|
|
|
#include "hip/hip_runtime_api.h"
|
|
#include "test_common.h"
|
|
|
|
#ifdef _WIN32
|
|
#define MAX_BURST_SIZE 40
|
|
#else
|
|
#define MAX_BURST_SIZE 100
|
|
#endif
|
|
|
|
// Create a lot of streams and then destroy 'em.
|
|
void createThenDestroyStreams(int iterations, int burstSize) {
|
|
hipStream_t* streams = new hipStream_t[burstSize];
|
|
|
|
for (int i = 0; i < iterations; i++) {
|
|
if (p_verbose & 0x1) {
|
|
printf("%s iter=%d, create %d then destroy %d\n", __func__, i, burstSize, burstSize);
|
|
}
|
|
for (int j = 0; j < burstSize; j++) {
|
|
if (p_verbose & 0x2) {
|
|
printf(" %d.%d streamCreate\n", i, j);
|
|
}
|
|
HIPCHECK(hipStreamCreate(&streams[j]));
|
|
}
|
|
for (int j = 0; j < burstSize; j++) {
|
|
if (p_verbose & 0x2) {
|
|
printf(" %d.%d streamDestroy\n", i, j);
|
|
}
|
|
HIPCHECK(hipStreamDestroy(streams[j]));
|
|
}
|
|
}
|
|
|
|
delete[] streams;
|
|
}
|
|
|
|
|
|
void waitStreams(int iterations) {
|
|
// Repeatedly sync and wait for all streams to complete.
|
|
// TO make this interesting, the test has other threads repeatedly adding and removing streams
|
|
// to the device.
|
|
for (int i = 0; i < iterations; i++) {
|
|
HIPCHECK(hipDeviceSynchronize());
|
|
}
|
|
}
|
|
|
|
|
|
// Create 3 streams, all creating and destroying streams on the same device.
|
|
// Some create many queue, some not many.
|
|
//
|
|
void multiThread_pyramid(bool serialize, int iters) {
|
|
printf("%s creating %d streams\n", __func__, iters * MAX_BURST_SIZE);
|
|
std::thread t1(createThenDestroyStreams, iters * 1, MAX_BURST_SIZE);
|
|
if (serialize) {
|
|
t1.join();
|
|
printf("t1 done\n");
|
|
}
|
|
|
|
std::thread t2(createThenDestroyStreams, iters * 10, 10);
|
|
if (serialize) {
|
|
t2.join();
|
|
printf("t2 done\n");
|
|
}
|
|
|
|
std::thread t3(createThenDestroyStreams, iters * 100, 1);
|
|
if (serialize) {
|
|
t3.join();
|
|
printf("t3 done\n");
|
|
}
|
|
|
|
if (!serialize) {
|
|
t1.join();
|
|
t2.join();
|
|
t3.join();
|
|
}
|
|
}
|
|
|
|
|
|
// Create 3 streams, all creating and destroying streams on the same device.
|
|
// Try to keep number of streams near zero, to cause problems.
|
|
void multiThread_nearzero(bool serialize, int iters) {
|
|
printf("%s creating %d streams x 3 threads\n", __func__, iters);
|
|
std::thread t1(createThenDestroyStreams, iters, 1);
|
|
if (serialize) {
|
|
t1.join();
|
|
printf("t1 done\n");
|
|
}
|
|
|
|
std::thread t2(createThenDestroyStreams, iters, 1);
|
|
if (serialize) {
|
|
t2.join();
|
|
printf("t2 done\n");
|
|
}
|
|
|
|
std::thread t3(waitStreams, iters * 50);
|
|
if (serialize) {
|
|
t3.join();
|
|
printf("t3 done\n");
|
|
}
|
|
|
|
if (!serialize) {
|
|
t1.join();
|
|
printf("t1 done\n");
|
|
t2.join();
|
|
printf("t2 done\n");
|
|
t3.join();
|
|
printf("t3 done\n");
|
|
}
|
|
}
|
|
|
|
int main(int argc, char* argv[]) {
|
|
HipTest::parseStandardArguments(argc, argv, true);
|
|
|
|
// Serial version, just call once:
|
|
if (p_tests & 0x1) {
|
|
printf("\ntest 0x1 : serial createThenDestroyStreams(10) \n");
|
|
createThenDestroyStreams(10, 10);
|
|
};
|
|
|
|
/*disable, this takess a while and if the next one works then no need to run serial*/
|
|
if (1 && (p_tests & 0x2)) {
|
|
printf("\ntest 0x2 : serialized multiThread_pyramid(1) \n");
|
|
multiThread_pyramid(true, 3);
|
|
}
|
|
|
|
if (p_tests & 0x4) {
|
|
printf("\ntest 0x4 : parallel multiThread_pyramid(1) \n");
|
|
multiThread_pyramid(false, 3);
|
|
}
|
|
|
|
// if (p_tests & 0x8) {
|
|
// printf ("test 0x8 : multiThread_pyramid(100) \n");
|
|
// multiThread_pyramid(false, 100);
|
|
// }
|
|
|
|
if (p_tests & 0x10) {
|
|
printf("\ntest 0x10 : parallel multiThread_nearzero(1000) \n");
|
|
multiThread_nearzero(false, 1000);
|
|
}
|
|
|
|
passed();
|
|
}
|