diff --git a/projects/hip/tests/performance/stream/hipPerfStreamConcurrency.cpp b/projects/hip/tests/performance/stream/hipPerfStreamConcurrency.cpp new file mode 100644 index 0000000000..81dc5626c7 --- /dev/null +++ b/projects/hip/tests/performance/stream/hipPerfStreamConcurrency.cpp @@ -0,0 +1,446 @@ +/* + Copyright (c) 2015-2020 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. + */ + +/* HIT_START + * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" +#include + + +typedef struct { + double x; + double y; + double width; +} coordRec; + +static coordRec coords[] = { + {0.0, 0.0, 0.00001}, // All black +}; + + +static unsigned int numCoords = sizeof(coords) / sizeof(coordRec); + +__global__ void mandelbrot(uint *out, uint width, float xPos, float yPos, + float xStep, float yStep, uint maxIter) { + + + int tid = (blockIdx.x * blockDim.x + threadIdx.x); + int i = tid % (width/4); + int j = tid / (width/4); + int4 veci = make_int4(4*i, 4*i+1, 4*i+2, 4*i+3); + int4 vecj = make_int4(j, j, j, j); + float4 x0; + x0.data[0] = (float)(xPos + xStep*veci.data[0]); + x0.data[1] = (float)(xPos + xStep*veci.data[1]); + x0.data[2] = (float)(xPos + xStep*veci.data[2]); + x0.data[3] = (float)(xPos + xStep*veci.data[3]); + float4 y0; + y0.data[0] = (float)(yPos + yStep*vecj.data[0]); + y0.data[1] = (float)(yPos + yStep*vecj.data[1]); + y0.data[2] = (float)(yPos + yStep*vecj.data[2]); + y0.data[3] = (float)(yPos + yStep*vecj.data[3]); + + float4 x = x0; + float4 y = y0; + + uint iter = 0; + float4 tmp; + int4 stay; + int4 ccount = make_int4(0, 0, 0, 0); + + float4 savx = x; + float4 savy = y; + + stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f); + stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f); + stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f); + stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f); + + for (iter = 0; (stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && (iter < maxIter); + iter+=16) { + + + x = savx; + y = savy; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + // Two iterations + tmp = x*x + x0 - y*y; + y = 2.0f * x * y + y0; + x = tmp*tmp + x0 - y*y; + y = 2.0f * tmp * y + y0; + + stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f); + stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f); + stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f); + stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f); + + savx.data[0] = (bool)(stay.data[0] ? x.data[0] : savx.data[0]); + savx.data[1] = (bool)(stay.data[1] ? x.data[1] : savx.data[1]); + savx.data[2] = (bool)(stay.data[2] ? x.data[2] : savx.data[2]); + savx.data[3] = (bool)(stay.data[3] ? x.data[3] : savx.data[3]); + + savy.data[0] = (bool)(stay.data[0] ? y.data[0] : savy.data[0]); + savy.data[1] = (bool)(stay.data[1] ? y.data[1] : savy.data[1]); + savy.data[2] = (bool)(stay.data[2] ? y.data[2] : savy.data[2]); + savy.data[3] = (bool)(stay.data[3] ? y.data[3] : savy.data[3]); + + ccount.data[0] -= stay.data[0]*16; + ccount.data[1] -= stay.data[1]*16; + ccount.data[2] -= stay.data[2]*16; + ccount.data[3] -= stay.data[3]*16; + } + + + // Handle remainder + if (!(stay.data[0] & stay.data[1] & stay.data[2] & stay.data[3])) + { + iter = 16; + do + { + x = savx; + y = savy; + stay.x = ((x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= 4.0f) && (ccount.data[0] < maxIter); + stay.y = ((x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= 4.0f) && (ccount.data[1] < maxIter); + stay.z = ((x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= 4.0f) && (ccount.data[2] < maxIter); + stay.w = ((x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= 4.0f) && (ccount.data[3] < maxIter); + tmp = x; + x = x*x + x0 - y*y; + y = 2.0f*tmp*y + y0; + ccount.data[0] += stay.data[0]; + ccount.data[1] += stay.data[1]; + ccount.data[2] += stay.data[2]; + ccount.data[3] += stay.data[3]; + iter--; + savx.data[0] = (stay.data[0] ? x.data[0] : savx.data[0]); + savx.data[1] = (stay.data[1] ? x.data[1] : savx.data[1]); + savx.data[2] = (stay.data[2] ? x.data[2] : savx.data[2]); + savx.data[3] = (stay.data[3] ? x.data[3] : savx.data[3]); + savy.data[0] = (stay.data[0] ? y.data[0] : savy.data[0]); + savy.data[1] = (stay.data[1] ? y.data[1] : savy.data[1]); + savy.data[2] = (stay.data[2] ? y.data[2] : savy.data[2]); + savy.data[3] = (stay.data[3] ? y.data[3] : savy.data[3]); + } while ((stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && iter); + } + + + uint4 *vecOut = (uint4 *)out; + + vecOut[tid].data[0] = (uint)(ccount.data[0]); + vecOut[tid].data[1] = (uint)(ccount.data[1]); + vecOut[tid].data[2] = (uint)(ccount.data[2]); + vecOut[tid].data[3] = (uint)(ccount.data[3]); +} + + +class hipPerfStreamConcurrency { + public: + hipPerfStreamConcurrency(); + ~hipPerfStreamConcurrency(); + + void setNumKernels(unsigned int num) { + numKernels = num; + } + void setNumStreams(unsigned int num) { + numStreams = num; + } + unsigned int getNumStreams() { + return numStreams; + } + + unsigned int getNumKernels() { + return numKernels; + } + + void open(int deviceID); + void run(unsigned int testCase, unsigned int deviceId); + void close(void); + + private: + void setData(void *ptr, unsigned int value); + void checkData(uint *ptr); + + unsigned int numKernels; + unsigned int numStreams; + + unsigned int width_; + unsigned int bufSize; + unsigned int maxIter; + unsigned int coordIdx; + unsigned long long totalIters; + int numCUs; + +}; + + +hipPerfStreamConcurrency::hipPerfStreamConcurrency() {} + +hipPerfStreamConcurrency::~hipPerfStreamConcurrency() {} + +void hipPerfStreamConcurrency::open(int deviceId) { + + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + std::cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return; + } + + + HIPCHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIPCHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId << std::endl; + + numCUs = props.multiProcessorCount; +} + + +void hipPerfStreamConcurrency::close() { +} + + +void hipPerfStreamConcurrency::run(unsigned int testCase,unsigned int deviceId) { + + int clkFrequency = 0; + unsigned int numStreams = getNumStreams(); + unsigned int numKernels = getNumKernels(); + + HIPCHECK(hipDeviceGetAttribute(&clkFrequency, hipDeviceAttributeClockRate, deviceId)); + + clkFrequency =(unsigned int)clkFrequency/1000; + + // Maximum iteration count + // maxIter = 8388608 * (engine_clock / 1000).serial execution + maxIter = (unsigned int)(((8388608 * ((float)clkFrequency / 1000)) * numCUs) / 128); + maxIter = (maxIter + 15) & ~15; + + hipStream_t streams[numStreams]; + + uint * hPtr[numKernels]; + uint * dPtr[numKernels]; + + // Width is divisible by 4 because the mandelbrot kernel processes 4 pixels at once. + width_ = 256; + + bufSize = width_ * sizeof(uint); + + // Create streams for concurrency + for (uint i = 0; i < numStreams; i++) { + HIPCHECK(hipStreamCreate(&streams[i])); + } + + + // Allocate memory on the host and device + for (uint i = 0; i < numKernels; i++) { + HIPCHECK(hipHostMalloc((void **)&hPtr[i], bufSize, hipHostMallocDefault)); + setData(hPtr[i], 0xdeadbeef); + HIPCHECK(hipMalloc((uint **)&dPtr[i], bufSize)) + } + + + // Prepare kernel launch parameters + int threads = (bufSize/sizeof(uint)); + int threads_per_block = 64; + int blocks = (threads/threads_per_block) + (threads % threads_per_block); + + coordIdx = testCase % numCoords; + float xStep = (float)(coords[coordIdx].width / (double)width_); + float yStep = (float)(-coords[coordIdx].width / (double)width_); + float xPos = (float)(coords[coordIdx].x - 0.5 * coords[coordIdx].width); + float yPos = (float)(coords[coordIdx].y + 0.5 * coords[coordIdx].width); + + // Copy memory asynchronously and concurrently from host to device + for (uint i = 0; i < numKernels; i++) { + HIPCHECK(hipMemcpyHtoDAsync(dPtr[i], hPtr[i], bufSize, streams[i % numStreams])); + } + + + // Synchronize to make sure all the copies are completed + for(uint i = 0; i < numStreams; i++) { + HIPCHECK(hipStreamSynchronize(streams[i])); + } + + // Warm-up kernel with lower iteration + if (testCase == 0) { + maxIter = 256; + } + + // Time the kernel execution + auto all_start = std::chrono::steady_clock::now(); + + for (uint i = 0; i < numKernels; i++) { + hipLaunchKernelGGL(mandelbrot, dim3(blocks), dim3(threads_per_block), 0, streams[i%numStreams], + dPtr[i], width_, xPos, yPos, xStep, yStep, maxIter); + } + + + // Synchronize all the concurrent streans to have completed execution + for(uint i = 0; i < numStreams; i++) { + HIPCHECK(hipStreamSynchronize(streams[i])); + } + + + auto all_end = std::chrono::steady_clock::now(); + std::chrono::duration all_kernel_time = all_end - all_start; + + // Copy data back from device to the host + for(uint i = 0; i < numKernels; i++) { + HIPCHECK(hipMemcpyDtoHAsync(hPtr[i] ,dPtr[i], bufSize, streams[i % numStreams])); + } + + + if (testCase != 0) { + std::cout <<"Measured time for " << numKernels <<" kernels (s) on " << numStreams <<" stream (s): " + << all_kernel_time.count() << std::endl; + } + + + unsigned long long expected = + (unsigned long long)width_ * (unsigned long long)maxIter; + + for(uint i = 0 ; i < numStreams; i++) { + HIPCHECK(hipStreamDestroy(streams[i])); + } + + + // Free host and device memory + for (uint i = 0; i < numKernels; i++) { + HIPCHECK(hipFree(hPtr[i])); + HIPCHECK(hipFree(dPtr[i])); + } + + +} + + +void hipPerfStreamConcurrency::setData(void *ptr, unsigned int value) { + unsigned int *ptr2 = (unsigned int *)ptr; + for (unsigned int i = 0; i < width_ ; i++) { + ptr2[i] = value; + } +} + + +void hipPerfStreamConcurrency::checkData(uint *ptr) { + totalIters = 0; + for (unsigned int i = 0; i < width_; i++) { + totalIters += ptr[i]; + } +} + + +int main(int argc, char* argv[]) { + hipPerfStreamConcurrency streamConcurrency; + int deviceId = 0; + + streamConcurrency.open(deviceId); + + for (unsigned int testCase = 0; testCase < 5; testCase++) { + + + switch (testCase) { + + + case 0: + // Warm-up kernel + streamConcurrency.setNumStreams(1); + streamConcurrency.setNumKernels(1); + break; + + case 1: + // default stream executes serially + streamConcurrency.setNumStreams(1); + streamConcurrency.setNumKernels(1); + break; + + case 2: + // 2-way concurrency + streamConcurrency.setNumStreams(2); + streamConcurrency.setNumKernels(2); + break; + + case 3: + // 4-way concurrency + streamConcurrency.setNumStreams(4); + streamConcurrency.setNumKernels(4); + break; + + case 4: + streamConcurrency.setNumStreams(2); + streamConcurrency.setNumKernels(4); + break; + + case 5: + break; + + default: + break; + } + streamConcurrency.run(testCase, deviceId); + + } + + + passed(); +} diff --git a/projects/hip/tests/performance/stream/hipPerfStreamCreateCopyDestroy.cpp b/projects/hip/tests/performance/stream/hipPerfStreamCreateCopyDestroy.cpp new file mode 100644 index 0000000000..bd194ba3fe --- /dev/null +++ b/projects/hip/tests/performance/stream/hipPerfStreamCreateCopyDestroy.cpp @@ -0,0 +1,133 @@ +/* + Copyright (c) 2015-2020 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. + */ + +/* HIT_START + * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +using namespace std; + +#define BufSize 0x1000 +#define Iterations 0x100 +#define TotalStreams 4 +#define TotalBufs 4 + + +class hipPerfStreamCreateCopyDestroy { + private: + unsigned int numBuffers_; + unsigned int numStreams_; + const size_t totalStreams_[TotalStreams]; + const size_t totalBuffers_[TotalBufs]; + public: + hipPerfStreamCreateCopyDestroy() : numBuffers_(0), numStreams_(0), + totalStreams_{1, 2, 4, 8}, + totalBuffers_{1, 100, 1000, 5000} {}; + ~hipPerfStreamCreateCopyDestroy() {}; + void open(int deviceID); + void run(unsigned int testNumber); +}; + +void hipPerfStreamCreateCopyDestroy::open(int deviceId) { + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + std::cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return; + } + + HIPCHECK(hipSetDevice(deviceId)); + hipDeviceProp_t props = {0}; + HIPCHECK(hipGetDeviceProperties(&props, deviceId)); + std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name + << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId << std::endl; +} + +void hipPerfStreamCreateCopyDestroy::run(unsigned int testNumber) { + numStreams_ = totalStreams_[testNumber % TotalStreams]; + size_t iter = Iterations / (numStreams_ * ((size_t)1 << (testNumber / TotalBufs + 1))); + hipStream_t streams[numStreams_]; + + numBuffers_ = totalBuffers_[testNumber / TotalBufs]; + float* dSrc[numBuffers_]; + size_t nBytes = BufSize * sizeof(float); + + for (size_t b = 0; b < numBuffers_; ++b) { + HIPCHECK(hipMalloc(&dSrc[b], nBytes)); + } + + float* hSrc; + hSrc = new float[nBytes]; + HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess); + for (size_t i = 0; i < BufSize; i++) { + hSrc[i] = 1.618f + i; + } + + auto start = std::chrono::steady_clock::now(); + + for (size_t i = 0; i < iter; ++i) { + for (size_t s = 0; s < numStreams_; ++s) { + HIPCHECK(hipStreamCreate(&streams[s])); + } + + for (size_t s = 0; s < numStreams_; ++s) { + for (size_t b = 0; b < numBuffers_; ++b) { + HIPCHECK(hipMemcpyWithStream(dSrc[b], hSrc, nBytes, hipMemcpyHostToDevice, streams[s])); + } + } + + for (size_t s = 0; s < numStreams_; ++s) { + HIPCHECK(hipStreamDestroy(streams[s])); + } + } + + auto end = std::chrono::steady_clock::now(); + std::chrono::duration diff = end - start; + + auto time = static_cast(diff.count() * 1000 / (iter * numStreams_)); + + cout << "Create+Copy+Destroy time for " << numStreams_ << " streams and " + << setw(4) << numBuffers_ << " buffers " << " and " << setw(4) + << iter << " iterations " << time << " (ms) " << endl; + + delete [] hSrc; + for (size_t b = 0; b < numBuffers_; ++b) { + HIPCHECK(hipFree(dSrc[b])); + } +} + +int main(int argc, char* argv[]) { + hipPerfStreamCreateCopyDestroy streamCCD; + + int deviceId = 0; + streamCCD.open(deviceId); + + for (auto testCase = 0; testCase < TotalStreams * TotalBufs; testCase++) { + streamCCD.run(testCase); + } + + passed(); +} diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp deleted file mode 100644 index 664378f233..0000000000 --- a/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp +++ /dev/null @@ -1,43 +0,0 @@ -/* -Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR -IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* - * Conformance test for checking functionality of - * hipError_t hipGetDevice(int *device); - */ - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include "test_common.h" - -int main() { - int numDevices = 0; - int device; - HIPCHECK(hipGetDeviceCount(&numDevices)); - for (int i = 0; i < numDevices; i++) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipGetDevice(&device)); - HIPASSERT(device == i); - } - passed(); -} diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp deleted file mode 100644 index 5ee086c8ac..0000000000 --- a/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp +++ /dev/null @@ -1,36 +0,0 @@ -/* -Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR -IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t EXCLUDE_HIP_PLATFORM - * HIT_END - */ - -#include "test_common.h" - -int main() { - int numDevices = 0; - HIPCHECK(hipGetDeviceCount(&numDevices)); - for (int i = 0; i < numDevices; i++) { - HIPCHECK(hipSetDevice(i)); - } - HIPASSERT(hipErrorInvalidDevice == hipSetDevice(numDevices)); - passed(); -} diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetGetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetGetDevice.cpp new file mode 100644 index 0000000000..4224c974b3 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/device/hipSetGetDevice.cpp @@ -0,0 +1,356 @@ +/* + * 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. + */ + +/* + * Test designed to run on Linux based platforms + * Verifies functionality of + * -- hipSetDevice and hipGetDevice with different ROCR_VISIBLE_DEVICES and + * HIP_VISIBLE_DEVICES values set + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST_NAMED: %t hipSetGetDevice-invalidDevice + * TEST_NAMED: %t hipSetGetDevice-allValidDevice + * TEST_NAMED: %t hipSetGetDevice-validDev1 --computeDevCnt 1 + * TEST_NAMED: %t hipSetGetDevice-validDev2 --computeDevCnt 2 + * TEST_NAMED: %t hipSetGetDevice-validDev3 --computeDevCnt 3 + * TEST_NAMED: %t hipSetGetDevice-validDev4 --computeDevCnt 4 + * TEST_NAMED: %t hipSetGetDevice-validDev5 --computeDevCnt 5 + * TEST_NAMED: %t hipSetGetDevice-validDev6 --computeDevCnt 6 + * TEST_NAMED: %t hipSetGetDevice-validDev7 --computeDevCnt 7 + * TEST_NAMED: %t hipSetGetDevice-validDev8 --computeDevCnt 8 + * HIT_END + */ + +#include "test_common.h" +#include +#include + +void getDeviceCount(int *numDevices) { + int fd[2], val = 0; + pipe(fd); + + pid_t childPid; + childPid = fork(); + + if (childPid > 0) { // parent + close(fd[1]); + read(fd[0], &val, sizeof(val)); + close(fd[0]); + *numDevices = val; + + } else if (childPid == 0) { // child + int devCnt = 0; + close(fd[0]); + + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + + hipGetDeviceCount(&devCnt); + + write(fd[1], &devCnt, sizeof(devCnt)); + close(fd[1]); + exit(0); + + } else { + failed("fork() failed. Exiting the test\n"); + } +} + +#define MAX_SIZE 30 + +// Pass either -1 in deviceNumber or invalid device number +bool testInvalidDevice(int numDevices, bool useRocrEnv, int deviceNumber) { + bool testResult = true; + int device; + int tempCount = 0; + int setDeviceErrorCheck = 0; + int getDeviceErrorCheck = 0; + int getDeviceCountErrorCheck = 0; + int fd[2]; + pipe(fd); + + pid_t cPid; + cPid = fork(); + + char visibleDeviceString[MAX_SIZE] = {}; + snprintf(visibleDeviceString, MAX_SIZE, "%d", deviceNumber); + + if (cPid == 0) { // child + hipError_t err; + + if (true == useRocrEnv) { + setenv("ROCR_VISIBLE_DEVICES", visibleDeviceString, 1); + } else { + setenv("HIP_VISIBLE_DEVICES", visibleDeviceString, 1); + } + + err = hipGetDeviceCount(&tempCount); + if (err == hipErrorNoDevice) { + getDeviceCountErrorCheck = 1; + } else { + printf("hipGetDeviceCount returns wrong value: %u\n", hipError_t(err)); + } + + for (int i = 0; i < numDevices; i++) { + err = hipSetDevice(i); + if (err == hipErrorInvalidDevice) { + setDeviceErrorCheck+= 1; + } else { + printf("hipSetDevice returns wrong value: %u\n", hipError_t(err)); + } + + err = hipGetDevice(&device); + if (err == hipErrorNoDevice) { + getDeviceErrorCheck+= 1; + } else { + printf("hipGetDevice returns wrong value: %u\n", hipError_t(err)); + } + } + + if ((getDeviceCountErrorCheck == 1) && (setDeviceErrorCheck == numDevices) + && (getDeviceErrorCheck == numDevices)) { + testResult = true; + + } else { + printf("Test failed for invalid device\n"); + testResult = false; + } + + close(fd[0]); + write(fd[1], &testResult, sizeof(testResult)); + close(fd[1]); + exit(0); + + } else if (cPid > 0) { // parent + close(fd[1]); + read(fd[0], &testResult, sizeof(testResult)); + close(fd[0]); + wait(NULL); + + } else { + printf("fork() failed\n"); + testResult = false; + } + return testResult; +} + +int deviceListLength = 1; +int parseExtraArguments(int argc, char* argv[]) { + int i = 0; + for (i = 1; i < argc; i++) { + const char* arg = argv[i]; + if (!strcmp(arg, " ")) { + // skip NULL args. + } else if (!strcmp(arg, "--computeDevCnt")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &deviceListLength)) { + failed("Bad deviceListLength argument"); + } + } else { + failed("Bad argument"); + } + } + return i; +} + +bool testValidDevices(int numDevices, bool useRocrEnv, int *deviceList, + int deviceListLength) { + bool testResult = true; + int tempCount = 0; + int device; + int setDeviceErrorCheck = 0; + int getDeviceErrorCheck = 0; + int getDeviceCountErrorCheck = 0; + int *deviceListPtr = deviceList; + char visibleDeviceString[MAX_SIZE] = {}; + + if ((NULL == deviceList) || ((deviceListLength < 1) || + deviceListLength > numDevices)) { + printf("Invalid argument for number of devices. Skipping current test\n"); + return testResult; + } + + for (int i = 0; i < deviceListLength; i++) { + if (NULL == deviceListPtr) { + printf("Invalid gpu index. Skipping current test\n"); + return testResult; + } + snprintf(visibleDeviceString + strlen(visibleDeviceString), MAX_SIZE, "%d,", + *deviceListPtr++); + } + + visibleDeviceString[strlen(visibleDeviceString)-1] = 0; + + int fd[2]; + pipe(fd); + + pid_t cPid; + cPid = fork(); + + if (cPid == 0) { + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + + if (true == useRocrEnv) { + setenv("ROCR_VISIBLE_DEVICES", visibleDeviceString, 1); + } else { + setenv("HIP_VISIBLE_DEVICES", visibleDeviceString, 1); + } + + hipError_t err; + err = hipGetDeviceCount(&tempCount); + + if (tempCount == deviceListLength) { + getDeviceCountErrorCheck = 1; + } else { + printf("hipGetDeviceCount failed. return value: %u\n", hipError_t(err)); + } + + for (int i = 0; i < numDevices; i++) { + err = hipSetDevice(i); + if (err == hipErrorInvalidDevice) { + setDeviceErrorCheck+= 1; + } + + err = hipGetDevice(&device); + if (err == hipErrorNoDevice) { + getDeviceErrorCheck+= 1; + } + } + + if ((getDeviceCountErrorCheck == 1) && (setDeviceErrorCheck == + (numDevices-deviceListLength)) && (getDeviceErrorCheck == 0)) { + testResult = true; + + } else { + printf("Test failed for device count %d\n", deviceListLength); + testResult = false; + } + + close(fd[0]); + write(fd[1], &testResult, sizeof(testResult)); + close(fd[1]); + exit(0); + + } else if (cPid > 0) { + close(fd[1]); + read(fd[0], &testResult, sizeof(testResult)); + close(fd[0]); + wait(NULL); + + } else { + printf("fork() failed\n"); + testResult = false; + } + return testResult; +} + +bool testValidDevicesBasic() { + bool testResult = true; + int numDevices = 0; + int device; + int validateCount = 0; + HIPCHECK(hipGetDeviceCount(&numDevices)); + printf("Available compute devices in the system: %d\n", numDevices); + + for (int i = 0; i < numDevices; i++) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipGetDevice(&device)); + if (device == i) { + validateCount+= 1; + } + } + if (numDevices != validateCount) { + testResult = false; + } + + return testResult; +} + +int main(int argc, char* argv[]) { + bool testResult = true; + int numDevices = 0; + int device; + int deviceList[MAX_SIZE]; + int extraArgs = 0; + +#ifdef __unix__ + getDeviceCount(&numDevices); + + if (numDevices == 0) { + failed("No gpus found. exiting\n"); + } + + printf("Available compute devices in the system: %d\n", numDevices); + + extraArgs = HipTest::parseStandardArguments(argc, argv, false); + parseExtraArguments(extraArgs, argv); + + if (extraArgs == 1) { + printf("\nRunning test for invalid compute device\n"); + // Test setting -1 to ROCR_VISIBLE_DEVICES + testResult &= testInvalidDevice(numDevices, true, -1); + // Test setting -1 to HIP_VISIBLE_DEVICES + testResult &= testInvalidDevice(numDevices, false, -1); + // Test setting invalid device to ROCR_VISIBLE_DEVICES + testResult &= testInvalidDevice(numDevices, true, numDevices); + // Test setting invalide device to HIP_VISIBLE_DEVICES + testResult &= testInvalidDevice(numDevices, false, numDevices); + + // Test for all available devices + printf("\nRunning test for all available compute devices\n"); + + for (int i = 0; i < numDevices; i++) { + deviceList[i] = i; + } + + testResult &= testValidDevices(numDevices, true, deviceList, numDevices); + testResult &= testValidDevices(numDevices, false, deviceList, numDevices); + } + + // Assigning values to deviceList in reverse order + + for (int i=0; i < deviceListLength; i++) { + deviceList[i] = deviceListLength-1-i; + } + + // Test for subset of available gpus + if (extraArgs == 3) { + printf("\nRunning test for %d compute devices\n", deviceListLength); + testResult &= testValidDevices(numDevices, true, deviceList, + deviceListLength); + testResult &= testValidDevices(numDevices, false, deviceList, + deviceListLength); + } + +#else + + printf("Running basic test on Windows\n"); + testResult &= testValidDevicesBasic(); + +#endif + + if (testResult == true) { + passed(); + + } else { + failed("One or more tests failed\n"); + } +} diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyNegativeMThrdMSize.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyNegativeMThrdMSize.cpp new file mode 100644 index 0000000000..1be5c3c3db --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyNegativeMThrdMSize.cpp @@ -0,0 +1,1188 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// Testcase Description: This test case achieves two scenarios +// 1) Verifies the working of Memcpy apis for range of Memory sizes from +// smallest one unit transfer to maxmem available. +// 2) Launches NUM_THREADS threads. Each thread in turn tests the working +// of 8 hipmemcpy apis + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_Negative_tests --tests 1 + * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_MultiThread_tests --tests 2 + * TEST_NAMED: %t hipMemcpyNegativeMThrdMSize_MultiSize_singleType --tests 3 --memcpyPeersOnly 0 --testAllTypes 0 + * HIT_END + */ + +#include +#include +#include +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define NUM_THREADS 10 +#define NUM_ELM 1024*1024 +#define HIPTEST_TRUE 1 + +int memcpyPeersOnly = 1; +int testAllTypes = 0; +int Available_Gpus = 0; +std::atomic failureCount{0}; + +enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D, + TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC, + TEST_MEMCPYD2DASYNC, TEST_MAX}; +std::vector apiNameToTest = { "hipMemcpy", "hipMemcpyH2D", + "hipMemcpyD2H", "hipMemcpyD2D", "hipMemcpyAsync", + "hipMemcpyH2DAsync", "hipMemcpyD2HAsync", "hipMemcpyD2DAsync" }; + +// If memcpyPeersOnly is true, then checks if given gpus are peers and returns +// true if they are peers, else false +// If memcpyPeersOnly is false, then returns true always +bool gpusIsPeer(int gpu0, int gpu1) { + bool bRet = true; + if (HIPTEST_TRUE == memcpyPeersOnly) { + int CanAccessPeer1 = 0, CanAccessPeer2 = 0; + HIPCHECK(hipDeviceCanAccessPeer(&CanAccessPeer1, gpu0, gpu1)); + HIPCHECK(hipDeviceCanAccessPeer(&CanAccessPeer2, gpu1, gpu0)); + if ((CanAccessPeer1 * CanAccessPeer2) == 0) { + bRet = false; + } + } + + return bRet; +} + +template +class memcpyTests { + public: + T *A_h, *B_h; + apiToTest api; + size_t NUM_ELMTS = 0; + hipStream_t stream; + memcpyTests(apiToTest val, size_t num_elmts); + bool Memcpy_And_verify(); + ~memcpyTests(); +}; + +class Memcpy_Negative_Tests { + public: + // The following function will test negative scenarios with hipMemcpy() + bool Test_Memcpy(void); + bool Test_MemcpyAsync(void); + bool Test_MemcpyHtoD(void); + bool Test_MemcpyHtoDAsync(void); + bool Test_MemcpyDtoH(void); + bool Test_MemcpyDtoHAsync(void); + bool Test_MemcpyDtoD(void); + bool Test_MemcpyDtoDAsync(void); +}; + + +bool Memcpy_Negative_Tests::Test_Memcpy(void) { + bool IfTestPassed = true; + std::string str_out, str_err = "hipErrorInvalidValue"; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + + if ((A_h == NULL) || (B_h == NULL)) { + failed("Malloc call failed!"); + } + + HIPCHECK(hipMalloc(&A_d, (NUM_ELM*sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM*sizeof(float)))); + + for ( int i = 0; i < NUM_ELM; ++i ) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Copying only half the memory on device side from host + HIPCHECK(hipMemcpy(A_d, A_h, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); + // Copying device memory to host to verify if the content is consistent + HIPCHECK(hipMemcpy(B_h, A_d, NUM_ELM * sizeof(float), hipMemcpyDefault)); + // Verifying the host content copied in the above step for consistency. + int Data_mismatch = 0; + + for (int i = 0; i < (NUM_ELM/2); ++i) { + if (B_h[i] != 123) { + Data_mismatch++; + } + } + + if (Data_mismatch != 0) { + printf("Data Mismatch for negative test\n"); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpy(NULL, A_d, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpy with NULL for destination parameter.\n"); + printf("Error: %s\n", str_out.c_str()); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpy(A_h, NULL, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpy with NULL for source\n"); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpy(NULL, NULL, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpy with NULL for source and destination\n"); + IfTestPassed = false; + } + + // To check the behaviour if both the ptrs provided are same + HIPCHECK(hipMemcpy(A_d, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); + HIPCHECK(hipMemcpy(A_h, A_h, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); + + // To check the consistency of the data + HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); + Data_mismatch = 0; + + for (int i = 0; i < (NUM_ELM/2); ++i) { + if (B_h[i] != 123) { + Data_mismatch++; + } + } + + if (Data_mismatch != 0) { + printf("Data Mismatch after memcpy of same src and destination\n"); + IfTestPassed = false; + } + + // Memory copy on same device with two different regions + HIPCHECK(hipMemcpy(A_d1, A_d, (NUM_ELM) * sizeof(float), hipMemcpyDefault)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyAsync(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + + HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Copying host data into the device. + HIPCHECK(hipMemcpyAsync(A_d1, A_h, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream)); + + // Passing null pointer: seg fault observed with the following. + str_out = hipGetErrorString(hipMemcpyAsync(NULL, A_h, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyAsync with NULL for destination\n"); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpyAsync(A_d, NULL, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyAsync with NULL for source\n"); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpyAsync(NULL, NULL, + NUM_ELM * sizeof(float), + hipMemcpyDefault, stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyAsync with NULL for source and destination\n"); + IfTestPassed = false; + } + + // Passing default stream just for sanity kind of check + HIPCHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM * sizeof(float), hipMemcpyDefault, + 0)); + + // Passing stream object belong to destination gpu + // which is against the suggested usage. + HIPCHECK(hipMemcpyAsync(A_d, A_d1, NUM_ELM * sizeof(float), + hipMemcpyDefault, stream)); + + // Passing incorrect memcpy kind is not allowed hence those scenarios + // are not included + + // Copying only half the memory on device side from host + HIPCHECK(hipMemcpyAsync(A_d, A_h, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault, stream)); + // Copying device memory to host to verify the content is consistent. + HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM/2) * sizeof(float), hipMemcpyDefault)); + + // Verifying the host content copied in the above step for consistency. + int Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM/2); ++i) { + if (B_h[i] != 123) { + Data_mismatch++; + } + } + + if (Data_mismatch != 0) { + printf("Data Mismatch after half the size memcpyAsync\n"); + IfTestPassed = false; + } + + // To check the behaviour if both the ptrs provided are same + HIPCHECK(hipMemcpyAsync(A_d, A_d, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault, stream)); + HIPCHECK(hipMemcpyAsync(A_h, A_h, (NUM_ELM/2) * sizeof(float), + hipMemcpyDefault, stream)); + // To check the consistency of the data + HIPCHECK(hipMemcpy(B_h, A_d, (NUM_ELM) * sizeof(float), hipMemcpyDefault)); + Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM); ++i) { + if (B_h[i] != 123) { + Data_mismatch++; + } + } + + if (Data_mismatch != 0) { + printf("Data Mismatch after memcpyAsync of same src and destination\n"); + IfTestPassed = false; + } + + // Memory copy on same device with two different regions + HIPCHECK(hipMemcpyAsync(A_d1, A_d, (NUM_ELM) * sizeof(float), + hipMemcpyDefault, stream)); + + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyHtoD(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Passing null ptr to check the API behavior. + // Expectation: It should not crash and exit gracefully. + str_out = hipGetErrorString(hipMemcpyHtoD(NULL, A_h, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyHtoD with NULL for destination\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyHtoD(A_d, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyHtoD with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyHtoD(NULL, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyHtoD with NULL for source and destination\n"); + IfTestPassed = false; + } + // Copy half of the allocated memory + HIPCHECK(hipMemcpyHtoD(A_d, A_h, NUM_ELM * sizeof(float) / 2)); + // copying back to host to verify + HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float))); + + int Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM / 2); ++i) + if (B_h[i] != 123) + Data_mismatch++; + + if (Data_mismatch != 0) { + printf("Data Mismatch after hipMemcpyHtoD with half size\n"); + IfTestPassed = false; + } + + // Setting device memory to zero + HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); + // Swap source and destination pointer + HIPCHECK(hipMemcpyHtoD(A_h, A_d, NUM_ELM * sizeof(float))); + + // Pass same pointers in source and destination params + HIPCHECK(hipMemcpyHtoD(A_h, A_h, NUM_ELM * sizeof(float))); + HIPCHECK(hipMemcpyHtoD(A_d, A_d, NUM_ELM * sizeof(float))); + + // Mem copy on same device with two different regions + HIPCHECK(hipMemcpyHtoD(A_d1, A_d, NUM_ELM * sizeof(float))); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyHtoDAsync(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Passing null ptr to check the API behavior. + // Expectation: It should not crash and exit gracefully. + str_out = hipGetErrorString(hipMemcpyHtoDAsync(NULL, A_h, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyHtoDAsync with NULL for destination\n"); + IfTestPassed = false; + } + + str_out = hipGetErrorString(hipMemcpyHtoDAsync(A_d, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyHtoDAsync with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyHtoDAsync(NULL, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed MemcpyHtoDAsync with NULL for source and destination\n"); + IfTestPassed = false; + } + + // Copy half of the allocated memory + HIPCHECK(hipMemcpyHtoDAsync(A_d, A_h, NUM_ELM * sizeof(float)/2, stream)); + // copying back to host to verify + HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float))); + int Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM/2); ++i) + if (B_h[i] != 123) + Data_mismatch++; + if (Data_mismatch != 0) { + printf("Data Mismatch after hipMemcpyHtoDAsync with half size\n"); + IfTestPassed = false; + } + + // Setting device memory to zero + HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); + // Swap source and destination pointer + HIPCHECK(hipMemcpyHtoDAsync(B_h, A_d, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipStreamSynchronize(stream)); + if (B_h[0] != 0) { + printf("Data Mismatch after hipMemcpyHtoDAsync with memset to 0\n"); + IfTestPassed = false; + } + + // Pass same pointers in source and destination params + HIPCHECK(hipMemcpyHtoDAsync(A_h, A_h, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipMemcpyHtoDAsync(A_d, A_d, NUM_ELM * sizeof(float), stream)); + + // Mem copy on same device with two different regions + HIPCHECK(hipMemcpyHtoDAsync(A_d1, A_d, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipStreamSynchronize(stream)); + // Checking the api with null stream + HIPCHECK(hipMemcpyHtoDAsync(A_d1, A_d, NUM_ELM * sizeof(float), 0)); + HIPCHECK(hipStreamSynchronize(stream)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyDtoH(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Copying data from host to device for further operations + HIPCHECK(hipMemcpyHtoD(A_d, A_h, NUM_ELM * sizeof(float))); + + // Passing null ptr to check the API behavior. + // Expectation: It should not crash and exit gracefully. + str_out = hipGetErrorString(hipMemcpyDtoH(NULL, A_d, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoH with NULL for destination\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoH(A_d, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoH with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoH(NULL, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoH with NULL for source and destination\n"); + IfTestPassed = false; + } + // Copy half of the allocated memory + HIPCHECK(hipMemcpyDtoH(B_h, A_d, NUM_ELM * sizeof(float)/2)); + + int Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM/2); ++i) + if (B_h[i] != 123) + Data_mismatch++; + + if (Data_mismatch != 0) { + printf("Data Mismatch after hipMemcpyDtoH with half size\n"); + IfTestPassed = false; + } + + // Setting device memory to zero + HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); + // Swap source and destination pointer + HIPCHECK(hipMemcpyDtoH(A_d, A_h, NUM_ELM * sizeof(float))); + + // Pass same pointers in source and destination params + HIPCHECK(hipMemcpyDtoH(A_h, A_h, NUM_ELM * sizeof(float))); + HIPCHECK(hipMemcpyDtoH(A_d, A_d, NUM_ELM * sizeof(float))); + + // Mem copy on same device with two diffeent regions + HIPCHECK(hipMemcpyDtoH(A_d1, A_d, NUM_ELM * sizeof(float))); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyDtoHAsync(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d = NULL, *A_d1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + + // Copying data from host to device for further operations + HIPCHECK(hipMemcpyHtoDAsync(A_d, A_h, NUM_ELM * sizeof(float), stream)); + + // Passing null ptr to check the API behavior. + // Expectation: It should not crash and exit gracefully. + str_out = hipGetErrorString(hipMemcpyDtoHAsync(NULL, A_d, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoHAsync with NULL for destination\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoHAsync(A_d, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoHAsync with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoHAsync(NULL, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed hipMemcpyDtoHAsync with NULL for source and destination\n"); + IfTestPassed = false; + } + + // Copy half of the allocated memory + HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d, NUM_ELM * sizeof(float)/2, stream)); + HIPCHECK(hipStreamSynchronize(stream)); + int Data_mismatch = 0; + for (int i = 0; i < (NUM_ELM/2); ++i) + if (B_h[i] != 123) + Data_mismatch++; + + if (Data_mismatch != 0) { + printf("Data Mismatch after hipMemcpyDtoHAsync with half size\n"); + IfTestPassed = false; + } + // Checking the api with default stream + HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d, NUM_ELM * sizeof(float), 0)); + // Setting device memory to zero + HIPCHECK(hipMemset(A_d, 0, NUM_ELM * sizeof(float))); + // Swap source and destination pointer + HIPCHECK(hipMemcpyDtoHAsync(A_d, A_h, NUM_ELM * sizeof(float), stream)); + + // Pass same pointers in source and destination params + HIPCHECK(hipMemcpyDtoHAsync(A_h, A_h, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipMemcpyDtoHAsync(A_d, A_d, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipStreamSynchronize(stream)); + // Mem copy on same device with two different regions + HIPCHECK(hipMemcpyDtoHAsync(A_d1, A_d, NUM_ELM * sizeof(float), stream)); + HIPCHECK(hipStreamSynchronize(stream)); + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(A_d1)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyDtoD(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d1 = NULL, *A_d2 = NULL, *Ad1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&Ad1, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMemset(A_d1, 0, NUM_ELM * sizeof(float))); + if (Available_Gpus > 1) { + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&A_d2, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMemset(A_d2, 1, NUM_ELM * sizeof(float))); + } + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + // Passing null pointers to check the behaviour:: + str_out = hipGetErrorString(hipMemcpyDtoD(&A_d1, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoD with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoD(NULL, &A_d2, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoD with NULL for destination\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoD(NULL, NULL, + NUM_ELM * sizeof(float))); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoD with NULL for source and destination\n"); + IfTestPassed = false; + } + + // Pass real but host ptr:: The below two scenarios gives seg fault. + // Behaviour is as expected + // HIPCHECK(hipMemcpyDtoD(&A_d1, &A_h, NUM_ELM * sizeof(float))); + // HIPCHECK(hipMemcpyDtoD(&A_h, &A_d1, NUM_ELM * sizeof(float))); + int Data_mismatch = 0; + // Copying half of actually allocated memory + HIPCHECK(hipSetDevice(0)); + if (Available_Gpus > 1) { + HIPCHECK(hipMemcpyHtoD(A_d1, A_h, NUM_ELM * sizeof(float))); + if (true == gpusIsPeer(0, 1)) { + HIPCHECK(hipMemcpyDtoD(A_d2, A_d1, NUM_ELM * sizeof(float)/2)); + HIPCHECK(hipMemcpyDtoH(B_h, A_d2, NUM_ELM * sizeof(float))); + for (int i = 0; i < NUM_ELM/2; ++i) { + if (B_h[i] != 123) + Data_mismatch++; + } + if (Data_mismatch != 0) { + printf("Data mismatch hipMemcpyDtoD between devices\n"); + IfTestPassed = false; + } + } + } + + // Passing same pointers for source and destination + HIPCHECK(hipMemcpyDtoD(A_d1, A_d1, NUM_ELM * sizeof(float))); + if (Available_Gpus > 1) { + HIPCHECK(hipMemcpyDtoD(A_d2, A_d2, NUM_ELM * sizeof(float))); + } + // Memcpy on same device with two different regions + HIPCHECK(hipMemcpyDtoD(Ad1, A_d1, NUM_ELM * sizeof(float))); + + HIPCHECK(hipFree(A_d1)); + HIPCHECK(hipFree(Ad1)); + if (Available_Gpus > 1) + HIPCHECK(hipFree(A_d2)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +bool Memcpy_Negative_Tests::Test_MemcpyDtoDAsync(void) { + bool IfTestPassed = true; + float *A_h = NULL, *B_h = NULL, *A_d1 = NULL, *A_d2 = NULL, *Ad1 = NULL; + std::string str_out, str_err = "hipErrorInvalidValue"; + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + A_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + B_h = reinterpret_cast(malloc(NUM_ELM * sizeof(float))); + if ((A_h == nullptr) || (B_h == nullptr)) { + failed("Malloc call failed!"); + } + HIPCHECK(hipMalloc(&A_d1, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMalloc(&Ad1, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMemset(A_d1, 0, NUM_ELM * sizeof(float))); + if (Available_Gpus > 1) { + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&A_d2, (NUM_ELM * sizeof(float)))); + HIPCHECK(hipMemset(A_d2, 1, NUM_ELM * sizeof(float))); + } + for (int i = 0; i < NUM_ELM; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } + // Passing null pointers to check the behaviour:: + str_out = hipGetErrorString(hipMemcpyDtoDAsync(&A_d1, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoDAsync with NULL for source\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoDAsync(NULL, &A_d2, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed for hipMemcpyDtoDAsync with NULL for destination\n"); + IfTestPassed = false; + } + str_out = hipGetErrorString(hipMemcpyDtoDAsync(NULL, NULL, + NUM_ELM * sizeof(float), + stream)); + if ((str_err.compare(str_out)) != 0) { + printf("Failed MemcpyDtoDAsync with NULL for source and destination\n"); + IfTestPassed = false; + } + + int Data_mismatch = 0; + // Copying half of actually allocated memory + HIPCHECK(hipSetDevice(0)); + if (Available_Gpus > 1) { + HIPCHECK(hipMemcpyHtoD(A_d1, A_h, NUM_ELM * sizeof(float))); + if (true == gpusIsPeer(0, 1)) { + HIPCHECK(hipMemcpyDtoDAsync(A_d2, A_d1, NUM_ELM * sizeof(float)/2, + stream)); + HIPCHECK(hipMemcpyDtoH(B_h, A_d2, NUM_ELM * sizeof(float))); + for (int i = 0; i < NUM_ELM/2; ++i) { + if (B_h[i] != 123) + Data_mismatch++; + } + if (Data_mismatch != 0) { + printf("Data mismatch hipMemcpyDtoDAsync between devices\n"); + IfTestPassed = false; + } + } + } + + // Memcpy on same device with two different regions + HIPCHECK(hipMemcpyDtoDAsync(Ad1, A_d1, NUM_ELM * sizeof(float) , stream)); + // Testing hipMemcpyDtoDAsync between two devices. + if (Available_Gpus > 1) { + if (true == gpusIsPeer(0, 1)) { + HIPCHECK(hipMemcpyDtoDAsync(A_d2, A_d1, NUM_ELM * sizeof(float), 0)); + } + } + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipFree(A_d1)); + HIPCHECK(hipFree(Ad1)); + if (Available_Gpus > 1) + HIPCHECK(hipFree(A_d2)); + free(A_h); + free(B_h); + + return IfTestPassed; +} + +template +memcpyTests::memcpyTests(apiToTest val, size_t num_elmts) { + api = val; + NUM_ELMTS = num_elmts; + printf("%zu ", NUM_ELMTS * sizeof(T)); + fflush(stdout); + A_h = reinterpret_cast(malloc(NUM_ELMTS * sizeof(T))); + B_h = reinterpret_cast(malloc(NUM_ELMTS * sizeof(T))); + if ((A_h == NULL) || (B_h == NULL)) { + exit(1); + } + + if (api >= TEST_MEMCPYD2D) { + HIPCHECK(hipStreamCreate(&stream)); + } + for (size_t i = 0; i < NUM_ELMTS; ++i) { + A_h[i] = 123; + B_h[i] = 0; + } +} + +template +bool memcpyTests::Memcpy_And_verify() { + bool bFail = false; + std::atomic Data_mismatch{0}; + T *A_d[Available_Gpus]; + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMalloc(&A_d[i], NUM_ELMTS * sizeof(T))); + } + HIPCHECK(hipSetDevice(0)); + + switch (api) { + case TEST_MEMCPY: // To test hipMemcpy() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + // HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost)); + for (int j = 0; j < NUM_ELMTS; ++j) { + if (A_h[j] != B_h[j]) { + Data_mismatch++; + } + } + + if (Data_mismatch.load() != 0) { + printf("hipMemcpy: Failed for GPU: %d\n", i); + bFail = true; + } + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = 1; j < Available_Gpus; ++j) { + if (true == gpusIsPeer(i, j)) { + HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDefault)); + // Copying in direction reverse of above to check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELMTS * sizeof(T), + hipMemcpyDefault)); + // Copying data to host to verify the content + HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELMTS * sizeof(T), + hipMemcpyDefault)); + for (int i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + + if (Data_mismatch.load() != 0) { + printf("hipMemcpy: Failed between GPU: %d and %d\n", i, j); + bFail = true; + } + } + } + } + break; + case TEST_MEMCPYH2D: // To test hipMemcpyHtoD() + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyHtoD(A_d[i], A_h, NUM_ELMTS * sizeof(T))); + // Copying data from device to host to check data consistency + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyHtoD: failed\n"); + bFail = true; + } + } + break; + case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpyDtoH(B_h, A_d[i], NUM_ELMTS * sizeof(T))); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyDtoH: failed\n"); + bFail = true; + } + } + break; + case TEST_MEMCPYD2D: // To test hipMemcpyDtoD() + if (Available_Gpus > 1) { + // First copy data from H to D and then from D to D followed by D to H + // HIPCHECK(hipMemcpyHtoD(A_d[0], A_h, NUM_ELMTS * sizeof(T))); + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = 1; j < Available_Gpus; ++j) { + if (true == gpusIsPeer(i, j)) { + HIPCHECK(hipMemcpyHtoD(A_d[i], A_h, NUM_ELMTS * sizeof(T))); + HIPCHECK(hipMemcpyDtoD(A_d[j], A_d[i], NUM_ELMTS * sizeof(T))); + // Copying in direction reverse of above to check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyDtoD(A_d[i], A_d[j], NUM_ELMTS * sizeof(T))); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyDtoD: failed between GPU: %d and %d\n", i, j); + bFail = true; + } + } + } + } + } else { + // As DtoD is not possible we will transfer data from HtH(A_h to B_h) + // so as to get through verification step + HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpy (Host to Host): failed\n"); + bFail = true; + } + } + break; + case TEST_MEMCPYASYNC: // To test hipMemcpyAsync() + // Copying data from host to individual devices followed by copying + // back to host and verifying the data consistency. + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyAsync(A_d[i], A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToDevice, stream)); + HIPCHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost, stream)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + + if (Data_mismatch.load() != 0) { + printf("hipMemcpyAsync: failed for GPU %d\n", i); + bFail = true; + } + } + // Device to Device copying for all combinations + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = 1; j < Available_Gpus; ++j) { + if (true == gpusIsPeer(i, j)) { + HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDefault, stream)); + // Copying in direction reverse of above to check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyAsync(A_d[i], A_d[j], NUM_ELMTS * sizeof(T), + hipMemcpyDefault, stream)); + HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELMTS * sizeof(T), + hipMemcpyDefault)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + + if (Data_mismatch.load() != 0) { + printf("hipMemcpyAsync: Failed between GPU: %d and %d\n", i, j); + bFail = true; + } + } + } + } + break; + case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync() + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpyHtoDAsync(A_d[i], A_h, NUM_ELMTS * sizeof(T), + stream)); + // Copying data from device to host to check data consistency + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyHtoDAsync: failed\n"); + bFail = true; + } + } + break; + case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync() + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpyDtoHAsync(B_h, A_d[i], NUM_ELMTS * sizeof(T), + stream)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyDtoHAsync: failed\n"); + bFail = true; + } + } + break; + case TEST_MEMCPYD2DASYNC: // To test hipMemcpyDtoDAsync() + if (Available_Gpus > 1) { + // First copy data from H to D and then from D to D followed by D to H + HIPCHECK(hipMemcpyHtoD(A_d[0], A_h, NUM_ELMTS * sizeof(T))); + for (int i = 0; i < Available_Gpus; ++i) { + for (int j = 1; j < Available_Gpus; ++j) { + if (true == gpusIsPeer(i, j)) { + HIPCHECK(hipSetDevice(j)); + HIPCHECK(hipMemcpyDtoDAsync(A_d[j], A_d[i], NUM_ELMTS * sizeof(T), + stream)); + // Copying in direction reverse of above to check if bidirectional + // access is happening without any error + HIPCHECK(hipMemcpyDtoDAsync(A_d[i], A_d[j], NUM_ELMTS * sizeof(T), + stream)); + HIPCHECK(hipDeviceSynchronize()); + HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELMTS * sizeof(T), + hipMemcpyDeviceToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpyDtoDAsync: failed GPU: %d and %d\n", i, j); + bFail = true; + } + } + } + } + } else { + // As DtoD is not possible we will transfer data from HtH(A_h to B_h) + // so as to get through verification step + HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELMTS * sizeof(T), + hipMemcpyHostToHost)); + for (size_t i = 0; i < NUM_ELMTS; ++i) { + if (A_h[i] != B_h[i]) + Data_mismatch++; + } + if (Data_mismatch.load() != 0) { + printf("hipMemcpy (Host to Host): failed\n"); + bFail = true; + } + } + break; + default: + printf("Did not receive valid option!\n"); + break; + } + + for (int i = 0; i < Available_Gpus; ++i) { + HIPCHECK(hipFree((A_d[i]))); + } + + // Return true if test is success + if (bFail == true) { + return false; + } else { + return true; + } +} + +template +memcpyTests::~memcpyTests() { + free(A_h); + free(B_h); + if (api >= TEST_MEMCPYD2D) { + HIPCHECK(hipStreamDestroy(stream)); + } +} + +void Thread_func(int Threadid) { + for (apiToTest api = TEST_MEMCPY; api < TEST_MAX; api = apiToTest(api + 1)) { + memcpyTests obj(api, 1024*1024); + if (false == obj.Memcpy_And_verify()) { + failureCount++; + } + } +} + +int parseExtraArguments(int argc, char* argv[]) { + int i = 0; + for (i = 1; i < argc; i++) { + const char* arg = argv[i]; + if (!strcmp(arg, " ")) { + // skip NULL args. + } else if (!strcmp(arg, "--memcpyPeersOnly")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &memcpyPeersOnly)) { + failed("Bad memcpyPeersOnly argument"); + } + } else if (!strcmp(arg, "--testAllTypes")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &testAllTypes)) { + failed("Bad testAllTypes argument"); + } + } else { + failed("Bad argument"); + } + } + return i; +} + + +int main(int argc, char* argv[]) { + bool TestPassed = true; + int extraArgs = 0; + HIPCHECK(hipGetDeviceCount(&Available_Gpus)); + extraArgs = HipTest::parseStandardArguments(argc, argv, false); + parseExtraArguments(extraArgs, argv); + + if (p_tests == 1) { + Memcpy_Negative_Tests test; + TestPassed = test.Test_Memcpy(); + TestPassed &= test.Test_MemcpyAsync(); + TestPassed &= test.Test_MemcpyHtoD(); + TestPassed &= test.Test_MemcpyHtoDAsync(); + TestPassed &= test.Test_MemcpyDtoD(); + TestPassed &= test.Test_MemcpyDtoDAsync(); + TestPassed &= test.Test_MemcpyDtoH(); + TestPassed &= test.Test_MemcpyDtoHAsync(); + if (TestPassed) { + passed(); + } else { + failed("Test Failed!"); + } + } else if (p_tests == 2) { + failureCount = 0; + std::thread Thrd[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) + Thrd[i] = std::thread(Thread_func, i); + + // Thread join is being called separately so as to allow the + // threads run parallely + for (int i = 0; i < NUM_THREADS; i++) + Thrd[i].join(); + + if (failureCount.load() != 0) { + failed("Failed"); + } else { + passed(); + } + } else if (p_tests == 3) { + size_t free = 0, total = 0; + HIPCHECK(hipMemGetInfo(&free, &total)); + failureCount = 0; + // Need to see if allocating all of available free memory will result in + // any issues in windows system before adding the same + std::vector NUM_ELMTS{1, 5, 10, 100, 1024, 10*1024, 100*1024, + 1024*1024, 10*1024*1024, 100*1024*1024, + 1024*1024*1024}; + + for (apiToTest api = TEST_MEMCPY; api < TEST_MAX; api = apiToTest(api+1)) { + printf("\nTesting %s for size: ", apiNameToTest[api].c_str()); + + // Check for 0 size + memcpyTests obj(api, 0); + obj.Memcpy_And_verify(); + HIPCHECK(hipDeviceSynchronize()); + + for (size_t x : NUM_ELMTS) { + if ((x * sizeof(char)) <= free) { + memcpyTests obj(api, x); + obj.Memcpy_And_verify(); + HIPCHECK(hipDeviceSynchronize()); + } + + if (HIPTEST_TRUE == testAllTypes) { + // Testing memcpy with various data types + if ((x * sizeof(int)) <= free) { + memcpyTests obj(api, x); + obj.Memcpy_And_verify(); + HIPCHECK(hipDeviceSynchronize()); + } + if ((x * sizeof(size_t)) <= free) { + memcpyTests obj(api, x); + obj.Memcpy_And_verify(); + HIPCHECK(hipDeviceSynchronize()); + } + if ((x * sizeof(long double)) <= free) { + memcpyTests obj(api, x); + obj.Memcpy_And_verify(); + HIPCHECK(hipDeviceSynchronize()); + } + } + } + } + printf("\n"); + passed(); + } else { + failed("Didnt receive any valid option\n"); + } +}