void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
bool syncBetweenCopies) {
size_t Nbytes = numElements * sizeof(T);
size_t eachCopyElements = numElements / numCopies;
size_t eachCopyBytes = eachCopyElements * sizeof(T);
printf(
"------------------------------------------------------------------------------------------"
"-----\n");
printf(
"testing: %s Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
__func__, Nbytes, (double)(Nbytes) / 1024.0 / 1024.0, numCopies, eachCopyElements,
eachCopyBytes);
T* A_d;
T *A_h1, *A_h2;
HIPCHECK(hipHostMalloc((void**)&A_h1, Nbytes, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&A_h2, Nbytes, hipHostMallocDefault));
HIPCHECK(hipMalloc(&A_d, Nbytes));
for (int i = 0; i < numElements; i++) {
A_h1[i] = 3.14f + static_cast<T>(i);
}
// stream=0; // fixme TODO
for (int i = 0; i < numCopies; i++) {
HIPASSERT(A_d + i * eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_d[i * eachCopyElements], &A_h1[i * eachCopyElements],
eachCopyBytes, hipMemcpyHostToDevice, stream));
}
if (syncBetweenCopies) {
HIPCHECK(hipDeviceSynchronize());
}
for (int i = 0; i < numCopies; i++) {
HIPASSERT(A_d + i * eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_h2[i * eachCopyElements], &A_d[i * eachCopyElements],
eachCopyBytes, hipMemcpyDeviceToHost, stream));
}
HIPCHECK(hipDeviceSynchronize());
// Verify we copied back all the data correctly:
for (int i = 0; i < numElements; i++) {
HIPASSERT(A_h1[i] == A_h2[i]);
}
HIPCHECK(hipHostFree(A_h1));
HIPCHECK(hipHostFree(A_h2));
HIPCHECK(hipFree(A_d));
}
void run(size_t size, hipStream_t stream1, hipStream_t stream2){
float *Ah, *Bh, *Cd, *Dd, *Eh;
float *Ahh, *Bhh, *Cdd, *Ddd, *Ehh;
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Ahh, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bhh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cdd, size));
HIPCHECK(hipMalloc(&Ddd, size));
HIPCHECK(hipHostMalloc((void**)&Ehh, size, hipHostMallocDefault));
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1));
HIPCHECK(hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream1, Cd);
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream2, Cdd);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1));
HIPCHECK(hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] = Ah[10] + 1.0f);
HIPASSERT(Ehh[10] = Ahh[10] + 1.0f);
}
int main() {
float *A, *Ad;
HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocDefault));
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
for (int i = 0; i < SIZE; i++) {
HIPCHECK(hipMemcpyAsync(Ad, A, SIZE, hipMemcpyHostToDevice, stream));
HIPCHECK(hipDeviceSynchronize());
}
}
void run1(size_t size, hipStream_t stream){
float *Ah, *Bh, *Cd, *Dd, *Eh;
HIPCHECK(hipHostMalloc((void**)&Ah, size, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&Bh, size, hipHostMallocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostMalloc((void**)&Eh, size, hipHostMallocDefault));
for(int i=0;i<N;i++){
Ah[i] = 1.0f;
}
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] == Ah[10] + 1.0f);
}
void simpleNegTest() {
printf("testing: %s\n", __func__);
hipError_t e;
float *A_malloc, *A_pinned, *A_d;
size_t Nbytes = N * sizeof(float);
A_malloc = (float*)malloc(Nbytes);
HIPCHECK(hipHostMalloc((void**)&A_pinned, Nbytes, hipHostMallocDefault));
A_d = NULL;
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPASSERT(A_d != NULL);
// Can't use default with async copy
e = hipMemcpyAsync(A_pinned, A_d, Nbytes, hipMemcpyDefault, NULL);
// HIPASSERT (e == hipSuccess);
// Not sure what happens here, the memory must be pinned.
e = hipMemcpyAsync(A_malloc, A_d, Nbytes, hipMemcpyDeviceToHost, NULL);
printf(" async memcpy of A_malloc to A_d. Result=%d\n", e);
// HIPASSERT (e==hipErrorInvalidValue);
}
void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs,
bool doHostSide) {
HIPASSERT(numElements % numInflight == 0); // Must be evenly divisible.
size_t Nbytes = numElements * sizeof(T);
size_t eachCopyElements = numElements / numInflight;
size_t eachCopyBytes = eachCopyElements * sizeof(T);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
printf(
"------------------------------------------------------------------------------------------"
"-----\n");
printf(
"testing: %s<%s> Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu "
"eachCopyBytes=%zu\n",
__func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes) / 1024.0 / 1024.0,
numPongs, numInflight, eachCopyElements, eachCopyBytes);
T* A_h = NULL;
T* A_d = NULL;
A_h = (T*)(HostTraits<AllocType>::Alloc(Nbytes));
HIPCHECK(hipMalloc(&A_d, Nbytes));
// Initialize the host array:
const T initValue = 13;
const T deviceConst = 2;
const T hostConst = 10000;
for (size_t i = 0; i < numElements; i++) {
A_h[i] = initValue + i;
}
for (int k = 0; k < numPongs; k++) {
for (int i = 0; i < numInflight; i++) {
HIPASSERT(A_d + i * eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_d[i * eachCopyElements], &A_h[i * eachCopyElements],
eachCopyBytes, hipMemcpyHostToDevice, stream));
}
hipLaunchKernel(addK<T>, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, 2,
numElements);
for (int i = 0; i < numInflight; i++) {
HIPASSERT(A_d + i * eachCopyElements < A_d + Nbytes);
HIPCHECK(hipMemcpyAsync(&A_h[i * eachCopyElements], &A_d[i * eachCopyElements],
eachCopyBytes, hipMemcpyDeviceToHost, stream));
}
if (doHostSide) {
assert(0);
#if 0
hipEvent_t e;
HIPCHECK(hipEventCreate(&e));
#endif
HIPCHECK(hipDeviceSynchronize());
for (size_t i = 0; i < numElements; i++) {
A_h[i] += hostConst;
}
}
};
HIPCHECK(hipDeviceSynchronize());
// Verify we copied back all the data correctly:
for (size_t i = 0; i < numElements; i++) {
T gold = initValue + i;
// Perform calcs in same order as test above to replicate FP order-of-operations:
for (int k = 0; k < numPongs; k++) {
gold += deviceConst;
if (doHostSide) {
gold += hostConst;
}
}
if (gold != A_h[i]) {
std::cout << i << ": gold=" << gold << " out=" << A_h[i] << std::endl;
HIPASSERT(gold == A_h[i]);
}
}
HIPCHECK(hipHostFree(A_h));
HIPCHECK(hipFree(A_d));
}
// IN: nStreams : number of streams to use for the test
// IN :useNullStream - use NULL stream. Synchronizes everything.
// IN: useSyncMemcpyH2D - use sync memcpy (no overlap) for H2D
// IN: useSyncMemcpyD2H - use sync memcpy (no overlap) for D2H
void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemcpyH2D,
bool useSyncMemcpyD2H) {
size_t Nbytes = N * sizeof(int);
printf("testing: %s(useNullStream=%d, useSyncMemcpyH2D=%d, useSyncMemcpyD2H=%d) ", __func__,
useNullStream, useSyncMemcpyH2D, useSyncMemcpyD2H);
printf("Nbytes=%zu (%6.1f MB)\n", Nbytes, (double)(Nbytes) / 1024.0 / 1024.0);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, true);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
hipStream_t* stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
if (useNullStream) {
nStreams = 1;
stream[0] = NULL;
} else {
for (int i = 0; i < nStreams; ++i) {
HIPCHECK(hipStreamCreate(&stream[i]));
}
}
size_t workLeft = N;
size_t workPerStream = N / nStreams;
for (int i = 0; i < nStreams; ++i) {
size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
size_t workBytes = work * sizeof(int);
size_t offset = i * workPerStream;
HIPASSERT(A_d + offset < A_d + Nbytes);
HIPASSERT(B_d + offset < B_d + Nbytes);
HIPASSERT(C_d + offset < C_d + Nbytes);
if (useSyncMemcpyH2D) {
HIPCHECK(hipMemcpy(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice));
} else {
HIPCHECK(hipMemcpyAsync(&A_d[offset], &A_h[offset], workBytes, hipMemcpyHostToDevice,
stream[i]));
HIPCHECK(hipMemcpyAsync(&B_d[offset], &B_h[offset], workBytes, hipMemcpyHostToDevice,
stream[i]));
};
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
&A_d[offset], &B_d[offset], &C_d[offset], work);
if (useSyncMemcpyD2H) {
HIPCHECK(hipMemcpy(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost));
} else {
HIPCHECK(hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost,
stream[i]));
}
}
HIPCHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, true);
free(stream);
};
int main() {
try {
bool done = false;
boost::fibers::fiber f1([&done]{
std::cout << "f1: entered" << std::endl;
try {
hipStream_t stream;
hipStreamCreate( & stream);
int size = 1024 * 1024;
int full_size = 20 * size;
int * host_a, * host_b, * host_c;
hipHostMalloc( & host_a, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_b, full_size * sizeof( int), hipHostMallocDefault);
hipHostMalloc( & host_c, full_size * sizeof( int), hipHostMallocDefault);
int * dev_a, * dev_b, * dev_c;
hipMalloc( & dev_a, size * sizeof( int) );
hipMalloc( & dev_b, size * sizeof( int) );
hipMalloc( & dev_c, size * sizeof( int) );
std::minstd_rand generator;
std::uniform_int_distribution<> distribution(1, 6);
for ( int i = 0; i < full_size; ++i) {
host_a[i] = distribution( generator);
host_b[i] = distribution( generator);
}
for ( int i = 0; i < full_size; i += size) {
hipMemcpyAsync( dev_a, host_a + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
hipMemcpyAsync( dev_b, host_b + i, size * sizeof( int), hipMemcpyHostToDevice, stream);
hipLaunchKernel( vector_add, dim3(size / 256), dim3(256), 0, stream, dev_a, dev_b, dev_c, size);
hipMemcpyAsync( host_c + i, dev_c, size * sizeof( int), hipMemcpyDeviceToHost, stream);
}
auto result = boost::fibers::hip::waitfor_all( stream);
BOOST_ASSERT( stream == std::get< 0 >( result) );
BOOST_ASSERT( hipSuccess == std::get< 1 >( result) );
std::cout << "f1: GPU computation finished" << std::endl;
hipHostFree( host_a);
hipHostFree( host_b);
hipHostFree( host_c);
hipFree( dev_a);
hipFree( dev_b);
hipFree( dev_c);
hipStreamDestroy( stream);
done = true;
} catch ( std::exception const& ex) {
std::cerr << "exception: " << ex.what() << std::endl;
}
std::cout << "f1: leaving" << std::endl;
});
boost::fibers::fiber f2([&done]{
std::cout << "f2: entered" << std::endl;
while ( ! done) {
std::cout << "f2: sleeping" << std::endl;
boost::this_fiber::sleep_for( std::chrono::milliseconds( 1 ) );
}
std::cout << "f2: leaving" << std::endl;
});
f1.join();
f2.join();
std::cout << "done." << std::endl;
return EXIT_SUCCESS;
} catch ( std::exception const& e) {
std::cerr << "exception: " << e.what() << std::endl;
} catch (...) {
std::cerr << "unhandled exception" << std::endl;
}
return EXIT_FAILURE;
}
