[ dtest ] Added different test scenario with hipMemcpyWithStream
- Test with one, two streams - Test Multi-gpu (one stream per gpu) - Test D-D (on same device/different devices). Can discover issues when devices are on same or different root complexes. - Test H-D/Default SWDEV-238517 for enhancing hip unit tests Change-Id: I8031a7eebe2f9c8c0e0996e2c7accb09ac0b96d4
このコミットが含まれているのは:
rohit pathania
Aryan Salmanpour
@@ -18,9 +18,9 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Conformance test for checking functionality of
|
||||
* hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t
|
||||
* sizeBytes);
|
||||
* Different test for checking functionality of
|
||||
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,hipMemcpyKind kind,
|
||||
* hipStream_t stream);
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
@@ -31,27 +31,475 @@ THE SOFTWARE.
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
int main() {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM);
|
||||
#define test_failed(test_name) printf("%s %s FAILED!%s\n", KRED, #test_name, KNRM);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice,stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
class HipMemcpyWithStreamtests {
|
||||
public:
|
||||
// Test hipMemcpyWithStream with one streams and launch kernel in
|
||||
// that stream, verify the data
|
||||
void TestwithOnestream(void);
|
||||
// Test hipMemcpyWithStream with two streams and launch kernels in
|
||||
// two streams, verify the data
|
||||
void TestwithTwoStream(void);
|
||||
// Test hipMemcpyWithStream with one stream for each gpu and launch
|
||||
// kernels in each, verify the data
|
||||
void TestOnMultiGPUwithOneStream(void);
|
||||
// Test hipMemcpyWithStream to copy data from device to host (hipMemcpyDeviceToHost)
|
||||
void TestkindDtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU
|
||||
void TestkindDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyHostToHost
|
||||
void TestkindHtoH(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault
|
||||
void TestkindDefault(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault for device to device transfer case
|
||||
void TestkindDefaultForDtoD(void);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device
|
||||
void TestDtoDonSameDevice(void);
|
||||
};
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
passed();
|
||||
void HipMemcpyWithStreamtests::TestwithOnestream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestwithTwoStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestDtoDonSameDevice(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int noOfstreams = 2;
|
||||
int *A_d[noOfstreams], *B_d[noOfstreams], *C_d[noOfstreams];
|
||||
int *A_h[noOfstreams], *B_h[noOfstreams], *C_h[noOfstreams];
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
|
||||
hipStream_t stream[noOfstreams];
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMalloc(&A_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[1], Nbytes));
|
||||
C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[1] != NULL);
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
|
||||
if (A_d[1]) {
|
||||
HIPCHECK(hipFree(A_d[1]));
|
||||
}
|
||||
if (B_d[1]) {
|
||||
HIPCHECK(hipFree(B_d[1]));
|
||||
}
|
||||
if (C_d[1]) {
|
||||
HIPCHECK(hipFree(C_d[1]));
|
||||
}
|
||||
if (C_h[1]) {
|
||||
free(C_h[1]);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < noOfstreams; ++i) {
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestOnMultiGPUwithOneStream(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i], &A_h[i], &B_h[i], &C_h[i], N, false);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i))
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDefault(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HIPCHECK(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
}
|
||||
|
||||
int *A_d[numDevices], *B_d[numDevices], *C_d[numDevices];
|
||||
int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
|
||||
|
||||
// Initialize and create the host and device elements for first device
|
||||
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
|
||||
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(C_h[i] != NULL);
|
||||
}
|
||||
|
||||
hipStream_t stream[numDevices];
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i]));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i]));
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
|
||||
static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
|
||||
}
|
||||
|
||||
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
|
||||
HIPCHECK(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamtests::TestkindHtoH(void) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_h, *B_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
|
||||
// Allocate memory to A_h and B_h
|
||||
A_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(A_h != NULL);
|
||||
B_h = static_cast<int*>(malloc(Nbytes));
|
||||
HIPASSERT(B_h != NULL);
|
||||
|
||||
for (size_t i = 0; i < N; ++i) {
|
||||
if (A_h) (A_h)[i] = 3.146f + i; // Pi
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
HIPASSERT(A_h[i] == B_h[i]);
|
||||
}
|
||||
|
||||
if (A_h) {
|
||||
free(A_h);
|
||||
}
|
||||
if (B_h) {
|
||||
free(B_h);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
int main() {
|
||||
HipMemcpyWithStreamtests tests;
|
||||
tests.TestwithOnestream();
|
||||
test_passed(TestwithOnestream);
|
||||
tests.TestwithTwoStream();
|
||||
test_passed(TestwithTwoStream);
|
||||
tests.TestkindDtoH();
|
||||
test_passed(TestkindsDtoH);
|
||||
tests.TestkindDefault();
|
||||
test_passed(TestkindDefault);
|
||||
tests.TestDtoDonSameDevice();
|
||||
test_passed(TestDtoDonSameDevice);
|
||||
tests.TestOnMultiGPUwithOneStream();
|
||||
test_passed(TestOnMultiGPUwithOneStream);
|
||||
tests.TestkindDtoD();
|
||||
test_passed(TestkindDtoD);
|
||||
tests.TestkindDefaultForDtoD();
|
||||
test_passed(TestkindDefaultForDtoD);
|
||||
tests.TestkindHtoH();
|
||||
test_passed(TestkindsHtoH);
|
||||
}
|
||||
|
||||
新しいイシューから参照
ユーザーをブロックする