agodavar
662 regels
21 KiB
C++
662 regels
21 KiB
C++
/*
|
|
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
|
Permission is hereby granted, free of charge, to any person obtaining a copy
|
|
of this software and associated documentation files (the "Software"), to deal
|
|
in the Software without restriction, including without limitation the rights
|
|
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
|
copies of the Software, and to permit persons to whom the Software is
|
|
furnished to do so, subject to the following conditions:
|
|
The above copyright notice and this permission notice shall be included in
|
|
all copies or substantial portions of the Software.
|
|
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
|
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
|
THE SOFTWARE.
|
|
*/
|
|
|
|
/*
|
|
* Different test for checking functionality of
|
|
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
|
|
* hipMemcpyKind kind, hipStream_t stream);
|
|
*/
|
|
|
|
/* HIT_START
|
|
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvidia
|
|
* TEST: %t
|
|
* HIT_END
|
|
*/
|
|
|
|
#include <vector>
|
|
#include <thread>
|
|
#include <chrono>
|
|
#include "test_common.h"
|
|
|
|
#define LEN 64
|
|
#define SIZE LEN << 2
|
|
#define THREADS 2
|
|
#define MAX_THREADS 16
|
|
|
|
|
|
#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);
|
|
|
|
enum class ops
|
|
{ TestwithOnestream,
|
|
TestwithTwoStream,
|
|
TestOnMultiGPUwithOneStream,
|
|
TestkindDtoH,
|
|
TestkindDtoD,
|
|
TestkindHtoH,
|
|
TestkindDefault,
|
|
TestkindDefaultForDtoD,
|
|
TestDtoDonSameDevice,
|
|
END_OF_LIST
|
|
};
|
|
|
|
|
|
class HipMemcpyWithStreamMultiThreadtests {
|
|
// 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);
|
|
|
|
public:
|
|
// run all the tests on multithreaded.
|
|
void TestwithMultiThreaded(ops op);
|
|
};
|
|
|
|
struct joinable_thread : std::thread {
|
|
template <class... Xs>
|
|
explicit joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...)
|
|
{} // NOLINT
|
|
|
|
joinable_thread& operator=(joinable_thread&& other) = default;
|
|
joinable_thread(joinable_thread&& other) = default;
|
|
|
|
~joinable_thread() {
|
|
if (this->joinable())
|
|
this->join();
|
|
}
|
|
};
|
|
|
|
void HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::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]));
|
|
|
|
// Copying device data from 1st GPU to the rest of the the GPUs that is
|
|
// numDevices in the setup. 1st GPU start numbering from 0,1,2..n etc.
|
|
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]));
|
|
}
|
|
|
|
|
|
// Launching the kernel including the 1st GPU to the no of GPUs present
|
|
// in the setup. 1st GPU start numbering from 0,1,2..n etc.
|
|
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 HipMemcpyWithStreamMultiThreadtests::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 HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) {
|
|
size_t Nbytes = N * sizeof(int);
|
|
int numDevices = 0;
|
|
|
|
|
|
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
|
HIPCHECK(hipGetDeviceCount(&numDevices));
|
|
// Test case will not run on single GPU setup.
|
|
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(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);
|
|
}
|
|
|
|
hipStream_t stream[numDevices];
|
|
for (int i=0; i < numDevices; ++i) {
|
|
HIPCHECK(hipSetDevice(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]));
|
|
|
|
// Copying device data from 1st GPU to the rest of the the GPUs
|
|
// using hipMemcpyDefault kind that is numDevices in the setup.
|
|
// 1st GPU start numbering from 0,1,2..n etc.
|
|
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));
|
|
// Output of each GPU is getting validated with input of 1st GPU.
|
|
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 HipMemcpyWithStreamMultiThreadtests::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));
|
|
}
|
|
|
|
|
|
void HipMemcpyWithStreamMultiThreadtests::TestwithMultiThreaded(ops op) {
|
|
int n = min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS);
|
|
std::vector<joinable_thread> threads;
|
|
|
|
for (uint32_t i = 0; i < n; i++) {
|
|
threads.emplace_back(std::thread{[&] {
|
|
switch ( op ) {
|
|
case ops::TestwithOnestream:
|
|
TestwithOnestream();
|
|
break;
|
|
case ops::TestwithTwoStream:
|
|
TestwithTwoStream();
|
|
break;
|
|
case ops::TestkindDtoH:
|
|
TestkindDtoH();
|
|
break;
|
|
case ops::TestkindHtoH:
|
|
TestkindHtoH();
|
|
break;
|
|
case ops::TestkindDtoD:
|
|
TestkindDtoD();
|
|
break;
|
|
case ops::TestOnMultiGPUwithOneStream:
|
|
TestOnMultiGPUwithOneStream();
|
|
break;
|
|
case ops::TestkindDefault:
|
|
TestkindDefault();
|
|
break;
|
|
case ops::TestkindDefaultForDtoD:
|
|
TestkindDefaultForDtoD();
|
|
break;
|
|
case ops::TestDtoDonSameDevice:
|
|
TestDtoDonSameDevice();
|
|
break;
|
|
default:{}
|
|
}
|
|
}});
|
|
}
|
|
}
|
|
|
|
|
|
int main() {
|
|
HipMemcpyWithStreamMultiThreadtests tests;
|
|
for (int op = static_cast<int>(ops::TestwithOnestream);
|
|
op < static_cast<int>(ops::END_OF_LIST); ++op) {
|
|
tests.TestwithMultiThreaded(static_cast<ops>(op));
|
|
switch ( static_cast<ops>(op) ) {
|
|
case ops::TestwithOnestream:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestwithOnestream);
|
|
break;
|
|
case ops::TestwithTwoStream:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestwithTwoStream);
|
|
break;
|
|
case ops::TestkindDtoH:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestkindDtoH);
|
|
break;
|
|
case ops::TestkindHtoH:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestkindHtoH);
|
|
break;
|
|
case ops::TestkindDtoD:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestkindDtoD);
|
|
break;
|
|
case ops::TestOnMultiGPUwithOneStream:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestOnMultiGPUwithOneStream);
|
|
break;
|
|
case ops::TestkindDefault:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestkindDefault);
|
|
break;
|
|
case ops::TestkindDefaultForDtoD:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestkindDefaultForDtoD);
|
|
break;
|
|
case ops::TestDtoDonSameDevice:
|
|
test_passed(HipMemcpyWithStreamMultiThreadtests
|
|
::TestDtoDonSameDevice);
|
|
break;
|
|
default: { test_failed("No Operation to done with API"); }
|
|
}
|
|
}
|
|
}
|