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 runTests(int64_t numElements)
{
size_t sizeBytes = numElements * sizeof(int);
printf ("\n\ntest: starting sequence with sizeBytes=%zu bytes, %6.2f MB\n", sizeBytes, sizeBytes/1024.0/1024.0);
int *C_h, *C_d;
HIPCHECK(hipMalloc(&C_d, sizeBytes));
HIPCHECK(hipHostMalloc(&C_h, sizeBytes));
{
test (0x01, C_d, C_h, numElements, syncNone, true /*expectMismatch*/);
test (0x02, C_d, C_h, numElements, syncNullStream, false /*expectMismatch*/);
test (0x04, C_d, C_h, numElements, syncOtherStream, true /*expectMismatch*/);
test (0x08, C_d, C_h, numElements, syncDevice, false /*expectMismatch*/);
// Sending a marker to to null stream may synchronize the otherStream
// - other created with hipStreamNonBlocking=0 : synchronization, should match
// - other created with hipStreamNonBlocking=1 : no synchronization, may mismatch
test (0x10, C_d, C_h, numElements, syncMarkerThenOtherStream, false /*expectMismatch*/);
// TODO - review why this test seems flaky
//test (0x20, C_d, C_h, numElements, syncMarkerThenOtherNonBlockingStream, true /*expectMismatch*/);
}
HIPCHECK(hipFree(C_d));
HIPCHECK(hipHostFree(C_h));
}
// TODO - deprecated function.
hipError_t hipFreeHost(void* ptr)
{
return hipHostFree(ptr);
}
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));
}
int main()
{
int *A, *Am, *B, *Ad, *C, *Cm;
A = new int[NUM];
B = new int[NUM];
C = new int[NUM];
for(int i=0;i<NUM;i++) {
A[i] = -1*i;
B[i] = 0;
C[i] = 0;
}
hipMalloc((void**)&Ad, SIZE);
hipHostMalloc((void**)&Am, SIZE);
hipHostMalloc((void**)&Cm, SIZE);
for(int i=0;i<NUM;i++) {
Am[i] = -1*i;
Cm[i] = 0;
}
hipStream_t stream;
hipStreamCreate(&stream);
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream);
hipStreamSynchronize(stream);
hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
hipStreamSynchronize(stream);
for(int i=0;i<NUM;i++) {
assert(Am[i] == B[i]);
assert(Am[i] == Cm[i]);
}
for(int i=0;i<NUM;i++) {
A[i] = -2*i;
B[i] = 0;
}
hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice);
hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
for(int i=0;i<NUM;i++) {
assert(A[i] == B[i]);
assert(A[i] == C[i]);
}
for(int i=0;i<NUM;i++) {
A[i] = -3*i;
B[i] = 0;
}
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream);
hipStreamSynchronize(stream);
hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
hipStreamSynchronize(stream);
for(int i=0;i<NUM;i++) {
assert(A[i] == B[i]);
assert(A[i] == C[i]);
}
hipHostFree(Am);
hipHostFree(Cm);
hipFree(Ad);
delete[] A;
delete[] B;
delete[] C;
passed();
}
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;
}
