From a8da9c1c3b2efec0a9cd06527d458cf5abe4cb7a Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Mon, 6 Nov 2017 11:31:52 +0530 Subject: [PATCH] More tests for memory,stream & error APIs [ROCm/hip commit: 093279b10bc3a5232c483539682d3cc03d1b0901] --- .../runtimeApi/error/hipPeekAtLastError.cpp | 38 +++++++++ .../runtimeApi/memory/hipMemcpyDtoDtoH.cpp | 77 ++++++++++++++++++ .../memory/hipMemcpyDtoDtoHAsync.cpp | 81 +++++++++++++++++++ .../src/runtimeApi/memory/hipMemcpyPeer.cpp | 77 ++++++++++++++++++ .../runtimeApi/memory/hipMemcpyPeerAsync.cpp | 79 ++++++++++++++++++ .../runtimeApi/stream/hipStreamGetFlags.cpp | 44 ++++++++++ 6 files changed, 396 insertions(+) create mode 100644 projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoH.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoHAsync.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp create mode 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp create mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp diff --git a/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp b/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp new file mode 100644 index 0000000000..f06ee3b271 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp @@ -0,0 +1,38 @@ +/* +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 + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +int main() +{ + hipSetDevice(-1); + if( hipPeekAtLastError() != hipSuccess) + passed(); +} diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoH.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoH.cpp new file mode 100644 index 0000000000..568ecb9c35 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoH.cpp @@ -0,0 +1,77 @@ +/* +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 hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes); + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +int main() +{ + hipDevice_t device; + size_t Nbytes = N*sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d; + int *A_h, *B_h, *C_h ; + + HIPCHECK(hipGetDeviceCount(&numDevices)); + if(numDevices > 1) + { + HIPCHECK(hipSetDevice(0)); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&X_d,Nbytes)); + HIPCHECK(hipMalloc(&Y_d,Nbytes)); + HIPCHECK(hipMalloc(&Z_d,Nbytes)); + + + HIPCHECK(hipSetDevice(0)); + HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes)); + HIPCHECK(hipMemcpyDtoD(Y_d, B_d, Nbytes)); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + HIPCHECK(hipMemcpyDtoH(C_h, Z_d, Nbytes)); + HIPCHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + } + + passed(); + +} + diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoHAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoHAsync.cpp new file mode 100644 index 0000000000..036ec16529 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDtoHAsync.cpp @@ -0,0 +1,81 @@ +/* +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 hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes); + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +int main() +{ + hipDevice_t device; + size_t Nbytes = N*sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d; + int *A_h, *B_h, *C_h ; + hipStream_t s; + + HIPCHECK(hipGetDeviceCount(&numDevices)); + if(numDevices > 1) + { + HIPCHECK(hipSetDevice(0)); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&X_d,Nbytes)); + HIPCHECK(hipMalloc(&Y_d,Nbytes)); + HIPCHECK(hipMalloc(&Z_d,Nbytes)); + + + HIPCHECK(hipSetDevice(0)); + HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HIPCHECK(hipStreamCreate(&s)); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s)); + HIPCHECK(hipMemcpyDtoDAsync(Y_d, B_d, Nbytes, s)); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + HIPCHECK(hipMemcpyDtoHAsync(C_h, Z_d, Nbytes, s)); + HIPCHECK(hipStreamSynchronize(s)); + HIPCHECK(hipDeviceSynchronize()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + } + + passed(); + + +} + diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp new file mode 100644 index 0000000000..dc2ba78ae0 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp @@ -0,0 +1,77 @@ +/* +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 hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes); + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +int main() +{ + hipDevice_t device; + size_t Nbytes = N*sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d; + int *A_h, *B_h, *C_h ; + + HIPCHECK(hipGetDeviceCount(&numDevices)); + if(numDevices > 1) + { + HIPCHECK(hipSetDevice(0)); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&X_d,Nbytes)); + HIPCHECK(hipMalloc(&Y_d,Nbytes)); + HIPCHECK(hipMalloc(&Z_d,Nbytes)); + + + HIPCHECK(hipSetDevice(0)); + HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + + HIPCHECK(hipSetDevice(1)); + hipMemcpyPeer(X_d, 1, A_d, 0, Nbytes); //this call is eqv to hipMemcpy(hipMemcpyD2D) which goes via stg bufs. + hipMemcpyPeer(Y_d, 1, B_d, 0, Nbytes); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + HIPCHECK ( hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + } + passed(); + + + + +} + diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp new file mode 100644 index 0000000000..1f0108c7cb --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp @@ -0,0 +1,79 @@ +/* +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 hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes); + */ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + +int main() +{ + hipDevice_t device; + size_t Nbytes = N*sizeof(int); + int numDevices = 0; + int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d; + int *A_h, *B_h, *C_h ; + hipStream_t s; + + + HIPCHECK(hipGetDeviceCount(&numDevices)); + if(numDevices > 1) + { + HIPCHECK(hipSetDevice(0)); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMalloc(&X_d,Nbytes)); + HIPCHECK(hipMalloc(&Y_d,Nbytes)); + HIPCHECK(hipMalloc(&Z_d,Nbytes)); + + + HIPCHECK(hipSetDevice(0)); + HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N); + HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HIPCHECK(hipStreamCreate(&s)); + HIPCHECK(hipSetDevice(1)); + HIPCHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, Nbytes, s)); + HIPCHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, Nbytes, s)); + + hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N); + HIPCHECK ( hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipDeviceSynchronize()); + HIPCHECK (hipStreamSynchronize(s)); + HipTest::checkVectorADD(A_h, B_h, C_h, N); + } + + passed(); + + +} + diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp new file mode 100644 index 0000000000..9212c70e7f --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp @@ -0,0 +1,44 @@ +/* +Copyright (c) 2015-2016 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 ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include "test_common.h" + + +int main(int argc, char *argv[]) +{ + hipStream_t stream; + unsigned int flags; + HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault)); + HIPCHECK(hipStreamGetFlags(stream, &flags)); + HIPASSERT(flags == 0); + HIPCHECK(hipStreamDestroy(stream)); + + HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIPCHECK(hipStreamGetFlags(stream, &flags)); + HIPASSERT(flags == 1); + HIPCHECK(hipStreamDestroy(stream)); + + passed(); +}