SWDEV-228443 - Enhancing hip unit tests for Memory Allocation APIs (#2616)
Adding new testcases for hipMalloc3D,hipMalloc3DArray, hipArrayCreate,hipMallocPitch and hipMallocArray APIs Change-Id: Ia2cc8865d605272995aaf703dd26954d11ded2ea
Этот коммит содержится в:
ROCm CI Service Account
коммит произвёл
GitHub
GitHub
родитель
79e0466f51
Коммит
b49e8e9fdf
@@ -71,7 +71,27 @@ THE SOFTWARE.
|
||||
printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \
|
||||
abort(); \
|
||||
}
|
||||
|
||||
#if HT_NVIDIA
|
||||
#define CTX_CREATE() \
|
||||
hipCtx_t context;\
|
||||
initHipCtx(&context);
|
||||
#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context));
|
||||
#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array));
|
||||
#define HIP_TEX_REFERENCE hipTexRef
|
||||
#define HIP_ARRAY hiparray
|
||||
static void initHipCtx(hipCtx_t *pcontext) {
|
||||
HIPCHECK(hipInit(0));
|
||||
hipDevice_t device;
|
||||
HIPCHECK(hipDeviceGet(&device, 0));
|
||||
HIPCHECK(hipCtxCreate(pcontext, 0, device));
|
||||
}
|
||||
#else
|
||||
#define CTX_CREATE()
|
||||
#define CTX_DESTROY()
|
||||
#define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array));
|
||||
#define HIP_TEX_REFERENCE textureReference*
|
||||
#define HIP_ARRAY hipArray*
|
||||
#endif
|
||||
|
||||
|
||||
// Utility Functions
|
||||
|
||||
@@ -75,6 +75,11 @@ set(TEST_SRC
|
||||
hipHostMalloc.cc
|
||||
hipMemcpy.cc
|
||||
hipMemcpyAsync.cc
|
||||
hipMallocPitch.cc
|
||||
hipMallocArray.cc
|
||||
hipMalloc3D.cc
|
||||
hipMalloc3DArray.cc
|
||||
hipArrayCreate.cc
|
||||
)
|
||||
else()
|
||||
set(TEST_SRC
|
||||
@@ -129,6 +134,11 @@ set(TEST_SRC
|
||||
hipHostMalloc.cc
|
||||
hipMemcpy.cc
|
||||
hipMemcpyAsync.cc
|
||||
hipMallocPitch.cc
|
||||
hipMallocArray.cc
|
||||
hipMalloc3D.cc
|
||||
hipMalloc3DArray.cc
|
||||
hipArrayCreate.cc
|
||||
)
|
||||
endif()
|
||||
|
||||
|
||||
@@ -0,0 +1,148 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
|
||||
/*
|
||||
hipArrayCreate API test scenarios
|
||||
1. Negative Scenarios
|
||||
2. Allocating Small and big chunk data
|
||||
3. Multithreaded scenario
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
static constexpr auto NUM_W{4};
|
||||
static constexpr auto BIGNUM_W{100};
|
||||
static constexpr auto NUM_H{4};
|
||||
static constexpr auto BIGNUM_H{100};
|
||||
static constexpr auto ARRAY_LOOP{100};
|
||||
|
||||
/*
|
||||
* This API verifies memory allocations for small and
|
||||
* bigger chunks of data.
|
||||
* Two scenarios are verified in this API
|
||||
* 1. SmallArray: Allocates NUM_W*NUM_H in a loop and
|
||||
* releases the memory and verifies the meminfo.
|
||||
* 2. BigArray: Allocates BIGNUM_W*BIGNUM_H in a loop and
|
||||
* releases the memory and verifies the meminfo
|
||||
*
|
||||
* In both cases, the memory info before allocation and
|
||||
* after releasing the memory should be the same.
|
||||
*
|
||||
*/
|
||||
|
||||
static void ArrayCreate_DiffSizes(int gpu) {
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
std::vector<size_t> array_size;
|
||||
array_size.push_back(NUM_W);
|
||||
array_size.push_back(BIGNUM_W);
|
||||
for (auto &size : array_size) {
|
||||
HIP_ARRAY array[ARRAY_LOOP];
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
HIP_ARRAY_DESCRIPTOR desc;
|
||||
desc.NumChannels = 1;
|
||||
if (size == NUM_W) {
|
||||
desc.Width = NUM_W;
|
||||
desc.Height = NUM_H;
|
||||
} else {
|
||||
desc.Width = BIGNUM_W;
|
||||
desc.Height = BIGNUM_H;
|
||||
}
|
||||
desc.Format = HIP_AD_FORMAT_FLOAT;
|
||||
HIP_CHECK(hipArrayCreate(&array[i], &desc));
|
||||
}
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
ARRAY_DESTROY(array[i]);
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((pavail != avail)) {
|
||||
HIPASSERT(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/*Thread function*/
|
||||
static void ArrayCreateThreadFunc(int gpu) {
|
||||
ArrayCreate_DiffSizes(gpu);
|
||||
}
|
||||
|
||||
/* This testcase verifies hipArrayCreate API for small and big chunks data*/
|
||||
TEST_CASE("Unit_hipArrayCreate_DiffSizes") {
|
||||
ArrayCreate_DiffSizes(0);
|
||||
}
|
||||
|
||||
|
||||
/* This testcase verifies the negative scenarios of
|
||||
* hipArrayCreate API
|
||||
*/
|
||||
TEST_CASE("Unit_hipArrayCreate_Negative") {
|
||||
HIP_ARRAY_DESCRIPTOR desc;
|
||||
HIP_ARRAY array;
|
||||
desc.Format = HIP_AD_FORMAT_FLOAT;
|
||||
desc.NumChannels = 1;
|
||||
desc.Width = NUM_W;
|
||||
desc.Height = NUM_H;
|
||||
#if HT_NVIDIA
|
||||
SECTION("NullPointer to Array") {
|
||||
REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("NullPointer to Channel Descriptor") {
|
||||
REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess);
|
||||
}
|
||||
#endif
|
||||
SECTION("Width 0 for Array Descriptor") {
|
||||
desc.Width = 0;
|
||||
REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Invalid NumChannels") {
|
||||
desc.NumChannels = 3;
|
||||
REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess);
|
||||
}
|
||||
}
|
||||
/*
|
||||
This testcase verifies the hipArrayCreate API in multithreaded
|
||||
scenario by launching threads in parallel on multiple GPUs
|
||||
and verifies the hipArrayCreate API with small and big chunks data
|
||||
*/
|
||||
TEST_CASE("Unit_hipArrayCreate_MultiThread") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt = 0;
|
||||
|
||||
devCnt = HipTest::getDeviceCount();
|
||||
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(ArrayCreateThreadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if (pavail != avail) {
|
||||
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,133 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
|
||||
/*
|
||||
hipMalloc3D API test scenarios
|
||||
1. Basic Functionality
|
||||
2. Negative Scenarios
|
||||
3. Allocating Small and big chunk data
|
||||
4. Multithreaded scenario
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
static constexpr auto SMALL_SIZE{4};
|
||||
static constexpr auto CHUNK_LOOP{100};
|
||||
static constexpr auto BIG_SIZE{100};
|
||||
/*
|
||||
This API verifies hipMalloc3D API by allocating memory in smaller chunks for
|
||||
CHUNK_LOOP iterations and checks for the memory leaks by get the memory
|
||||
info before and after the hipMalloc3D API and the difference should
|
||||
match with the allocated memory
|
||||
*/
|
||||
static void MemoryAlloc3DDiffSizes(int gpu) {
|
||||
HIPCHECK(hipSetDevice(gpu));
|
||||
std::vector<size_t> array_size;
|
||||
array_size.push_back(SMALL_SIZE);
|
||||
array_size.push_back(BIG_SIZE);
|
||||
for (auto &sizes : array_size) {
|
||||
size_t width = sizes * sizeof(float);
|
||||
size_t height{sizes}, depth{sizes};
|
||||
hipPitchedPtr devPitchedPtr[CHUNK_LOOP];
|
||||
hipExtent extent = make_hipExtent(width, height, depth);
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < CHUNK_LOOP; i++) {
|
||||
HIPCHECK(hipMalloc3D(&devPitchedPtr[i], extent));
|
||||
}
|
||||
for (int i = 0; i < CHUNK_LOOP; i++) {
|
||||
HIPCHECK(hipFree(devPitchedPtr[i].ptr));
|
||||
}
|
||||
HIPCHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((pavail != avail)) {
|
||||
HIPASSERT(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static void Malloc3DThreadFunc(int gpu) {
|
||||
MemoryAlloc3DDiffSizes(gpu);
|
||||
}
|
||||
|
||||
/*
|
||||
* This verifies the negative scenarios of hipMalloc3D API
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3D_Negative") {
|
||||
size_t width = SMALL_SIZE * sizeof(char);
|
||||
size_t height{SMALL_SIZE}, depth{SMALL_SIZE};
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
|
||||
SECTION("Passing nullptr to device pitched pointer") {
|
||||
hipExtent extent = make_hipExtent(width, height, depth);
|
||||
REQUIRE(hipMalloc3D(nullptr, extent) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Passing Max values to extent") {
|
||||
hipExtent extent = make_hipExtent(std::numeric_limits<size_t>::max(),
|
||||
std::numeric_limits<size_t>::max(),
|
||||
std::numeric_limits<size_t>::max());
|
||||
REQUIRE(hipMalloc3D(&devPitchedPtr, extent) != hipSuccess);
|
||||
}
|
||||
}
|
||||
/*
|
||||
* This verifies the hipMalloc3D API by
|
||||
* assigning width,height and depth as 10
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3D_Basic") {
|
||||
size_t width = SMALL_SIZE * sizeof(char);
|
||||
size_t height{SMALL_SIZE}, depth{SMALL_SIZE};
|
||||
hipPitchedPtr devPitchedPtr;
|
||||
hipExtent extent = make_hipExtent(width, height, depth);
|
||||
REQUIRE(hipMalloc3D(&devPitchedPtr, extent) == hipSuccess);
|
||||
}
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMalloc3D API by allocating
|
||||
smaller and big chunk data.
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3D_SmallandBigChunks") {
|
||||
MemoryAlloc3DDiffSizes(0);
|
||||
}
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMalloc3D API in multithreaded
|
||||
scenario by launching threads in parallel on multiple GPUs
|
||||
and verifies the hipMalloc3D API with small and big chunks data
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3D_MultiThread") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt = 0;
|
||||
|
||||
devCnt = HipTest::getDeviceCount();
|
||||
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(Malloc3DThreadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if (pavail != avail) {
|
||||
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,195 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
|
||||
/*
|
||||
hipMalloc3DArray API test scenarios
|
||||
1. Basic Functionality
|
||||
2. Negative Scenarios
|
||||
3. Allocating Small and big chunk data
|
||||
4. Multithreaded scenario
|
||||
*/
|
||||
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
static constexpr auto ARRAY_SIZE{4};
|
||||
static constexpr auto BIG_ARRAY_SIZE{100};
|
||||
static constexpr auto ARRAY_LOOP{100};
|
||||
|
||||
/*
|
||||
* This API verifies memory allocations for small and
|
||||
* bigger chunks of data.
|
||||
* Two scenarios are verified in this API
|
||||
* 1. SmallArray: Allocates ARRAY_SIZE in a loop and
|
||||
* releases the memory and verifies the meminfo.
|
||||
* 2. BigArray: Allocates BIG_ARRAY_SIZE in a loop and
|
||||
* releases the memory and verifies the meminfo
|
||||
*
|
||||
* In both cases, the memory info before allocation and
|
||||
* after releasing the memory should be the same
|
||||
*
|
||||
*/
|
||||
static void Malloc3DArray_DiffSizes(int gpu) {
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
std::vector<int> array_size;
|
||||
array_size.push_back(ARRAY_SIZE);
|
||||
array_size.push_back(BIG_ARRAY_SIZE);
|
||||
for (auto &size : array_size) {
|
||||
int width{size}, height{size}, depth{size};
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
|
||||
0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *arr[ARRAY_LOOP];
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
HIP_CHECK(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width,
|
||||
height, depth), hipArrayDefault));
|
||||
}
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
HIP_CHECK(hipFreeArray(arr[i]));
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((pavail != avail)) {
|
||||
HIPASSERT(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/* Thread Function */
|
||||
static void Malloc3DArrayThreadFunc(int gpu) {
|
||||
Malloc3DArray_DiffSizes(gpu);
|
||||
}
|
||||
|
||||
/*
|
||||
* Verifies the negative scenarios of hipMalloc3DArray API
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3DArray_Negative") {
|
||||
constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE};
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
|
||||
0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *arr;
|
||||
#if HT_NVIDIA
|
||||
SECTION("NullPointer to Array") {
|
||||
REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width,
|
||||
height, depth), hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("NullPointer to Channel Descriptor") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width,
|
||||
height, depth), hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
#endif
|
||||
SECTION("Width 0 in hipExtent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0,
|
||||
height, width), hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Height 0 in hipExtent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
|
||||
0, width), hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Invalid Flag") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
|
||||
height, depth), 100) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Width,Height & Depth 0 in hipExtent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0,
|
||||
0, 0), hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Max int values to extent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc,
|
||||
make_hipExtent(std::numeric_limits<int>::max(),
|
||||
std::numeric_limits<int>::max(),
|
||||
std::numeric_limits<int>::max()),
|
||||
hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
}
|
||||
/*
|
||||
* Verifies the extent validation scenarios
|
||||
* 1. Passing depth as 0 would create 2D array
|
||||
* 2. Passing height and depth as 0 would create 1D array
|
||||
* from hipMalloc3DArray API
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3DArray_ExtentValidation") {
|
||||
constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE};
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
|
||||
0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *arr;
|
||||
|
||||
SECTION("Depth 0 in hipExtent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
|
||||
height, 0), hipArrayDefault) == hipSuccess);
|
||||
HIP_CHECK(hipFreeArray(arr));
|
||||
}
|
||||
|
||||
SECTION("Height & Depth 0 in hipExtent") {
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
|
||||
0, 0), hipArrayDefault) == hipSuccess);
|
||||
HIP_CHECK(hipFreeArray(arr));
|
||||
}
|
||||
}
|
||||
|
||||
/*
|
||||
* Verifies hipMalloc3DArray API by passing width,height
|
||||
* and depth as 10
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3DArray_Basic") {
|
||||
constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE};
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
|
||||
0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *arr;
|
||||
|
||||
REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width,
|
||||
height, depth), hipArrayDefault) == hipSuccess);
|
||||
HIP_CHECK(hipFreeArray(arr));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") {
|
||||
Malloc3DArray_DiffSizes(0);
|
||||
}
|
||||
/*
|
||||
This testcase verifies the hipMalloc3DArray API in multithreaded
|
||||
scenario by launching threads in parallel on multiple GPUs
|
||||
and verifies the hipMalloc3DArray API with small and big chunks data
|
||||
*/
|
||||
TEST_CASE("Unit_hipMalloc3DArray_MultiThread") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt = 0;
|
||||
devCnt = HipTest::getDeviceCount();
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if (pavail != avail) {
|
||||
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,172 @@
|
||||
/*
|
||||
Copyright (c) 2022 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.
|
||||
*/
|
||||
|
||||
/*
|
||||
hipMallocArray API test scenarios
|
||||
1. Basic Functionality
|
||||
2. Negative Scenarios
|
||||
3. Allocating Small and big chunk data
|
||||
4. Multithreaded scenario
|
||||
*/
|
||||
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
static constexpr auto NUM_W{4};
|
||||
static constexpr auto BIGNUM_W{100};
|
||||
static constexpr auto BIGNUM_H{100};
|
||||
static constexpr auto NUM_H{4};
|
||||
static constexpr auto ARRAY_LOOP{100};
|
||||
|
||||
/*
|
||||
* This API verifies memory allocations for small and
|
||||
* bigger chunks of data.
|
||||
* Two scenarios are verified in this API
|
||||
* 1. NUM_W(small Data): Allocates NUM_W*NUM_H in a loop and
|
||||
* releases the memory and verifies the meminfo.
|
||||
* 2. BIGNUM_W(big data): Allocates BIGNUM_W*BIGNUM_H in a loop and
|
||||
* releases the memory and verifies the meminfo
|
||||
*
|
||||
* In both cases, the memory info before allocation and
|
||||
* after releasing the memory should be the same
|
||||
*
|
||||
*/
|
||||
|
||||
static void MallocArray_DiffSizes(int gpu) {
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
std::vector<size_t> array_size;
|
||||
array_size.push_back(NUM_W);
|
||||
array_size.push_back(BIGNUM_W);
|
||||
for (auto &size : array_size) {
|
||||
hipArray* A_d[ARRAY_LOOP];
|
||||
size_t tot, avail, ptot, pavail;
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
if (size == NUM_W) {
|
||||
HIP_CHECK(hipMallocArray(&A_d[i], &desc,
|
||||
NUM_W, NUM_H,
|
||||
hipArrayDefault));
|
||||
} else {
|
||||
HIP_CHECK(hipMallocArray(&A_d[i], &desc,
|
||||
BIGNUM_W, BIGNUM_H,
|
||||
hipArrayDefault));
|
||||
}
|
||||
}
|
||||
for (int i = 0; i < ARRAY_LOOP; i++) {
|
||||
HIP_CHECK(hipFreeArray(A_d[i]));
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if ((pavail != avail)) {
|
||||
HIPASSERT(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static void MallocArrayThreadFunc(int gpu) {
|
||||
MallocArray_DiffSizes(gpu);
|
||||
}
|
||||
|
||||
/*
|
||||
* This testcase verifies the negative scenarios of
|
||||
* hipMallocArray API
|
||||
*/
|
||||
TEST_CASE("Unit_hipMallocArray_Negative") {
|
||||
hipArray* A_d;
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
#if HT_NVIDIA
|
||||
SECTION("NullPointer to Array") {
|
||||
REQUIRE(hipMallocArray(nullptr, &desc,
|
||||
NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("NullPointer to Channel Descriptor") {
|
||||
REQUIRE(hipMallocArray(&A_d, nullptr,
|
||||
NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
#endif
|
||||
SECTION("Width 0 in hipMallocArray") {
|
||||
REQUIRE(hipMallocArray(&A_d, &desc,
|
||||
0, NUM_H, hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Height 0 in hipMallocArray") {
|
||||
REQUIRE(hipMallocArray(&A_d, &desc,
|
||||
NUM_W, 0, hipArrayDefault) == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Invalid Flag") {
|
||||
REQUIRE(hipMallocArray(&A_d, &desc,
|
||||
NUM_W, NUM_H, 100) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Max int values") {
|
||||
REQUIRE(hipMallocArray(&A_d, &desc,
|
||||
std::numeric_limits<int>::max(),
|
||||
std::numeric_limits<int>::max(),
|
||||
hipArrayDefault) != hipSuccess);
|
||||
}
|
||||
}
|
||||
/*
|
||||
* This testcase verifies the basic scenario of
|
||||
* hipMallocArray API for different datatypes
|
||||
* of size 10
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipMallocArray_Basic",
|
||||
"", int, unsigned int, float) {
|
||||
hipArray* A_d;
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
|
||||
REQUIRE(hipMallocArray(&A_d, &desc,
|
||||
NUM_W, NUM_H,
|
||||
hipArrayDefault) == hipSuccess);
|
||||
HIP_CHECK(hipFreeArray(A_d));
|
||||
}
|
||||
|
||||
|
||||
TEST_CASE("Unit_hipMallocArray_DiffSizes") {
|
||||
MallocArray_DiffSizes(0);
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMallocArray API in multithreaded
|
||||
scenario by launching threads in parallel on multiple GPUs
|
||||
and verifies the hipMallocArray API with small and big chunks data
|
||||
*/
|
||||
TEST_CASE("Unit_hipMallocArray_MultiThread") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt = 0;
|
||||
devCnt = HipTest::getDeviceCount();
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(MallocArrayThreadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if (pavail != avail) {
|
||||
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,296 @@
|
||||
/*
|
||||
Copyright (c) 2022 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 Scenarios of hipMallocPitch API
|
||||
1. Negative Scenarios
|
||||
2. Basic Functionality Scenario
|
||||
3. Allocate memory using hipMallocPitch API, Launch Kernel validate result.
|
||||
4. Allocate Memory in small chunks and large chunks and check for possible memory leaks
|
||||
5. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables.
|
||||
6. Multithreaded scenario
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
|
||||
static constexpr auto SMALLCHUNK_NUMW{4};
|
||||
static constexpr auto SMALLCHUNK_NUMH{4};
|
||||
static constexpr auto LARGECHUNK_NUMW{1025};
|
||||
static constexpr auto LARGECHUNK_NUMH{1000};
|
||||
static constexpr auto NUM_W{10};
|
||||
static constexpr auto NUM_H{10};
|
||||
static constexpr auto COLUMNS{8};
|
||||
static constexpr auto ROWS{8};
|
||||
static constexpr auto CHUNK_LOOP{100};
|
||||
|
||||
|
||||
template<typename T>
|
||||
__global__ void copy_var(T* A, T* B,
|
||||
size_t ROWS, size_t pitch_A) {
|
||||
for (uint64_t i = 0; i< ROWS*pitch_A; i= i+pitch_A) {
|
||||
A[i] = B[i];
|
||||
}
|
||||
}
|
||||
template<typename T>
|
||||
static bool validateResult(T* A, T* B, size_t pitch_A) {
|
||||
bool testResult = true;
|
||||
for (uint64_t i=0; i < pitch_A*ROWS; i=i+pitch_A) {
|
||||
if (A[i] != B[i]) {
|
||||
testResult = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
return testResult;
|
||||
}
|
||||
/*
|
||||
* This API verifies memory allocations for small and
|
||||
* bigger chunks of data.
|
||||
* Two scenarios are verified in this API
|
||||
* 1. SmallChunk: Allocates SMALLCHUNK_NUMW in a loop and
|
||||
* releases the memory and verifies the meminfo.
|
||||
* 2. LargeChunk: Allocates LARGECHUNK_NUMW in a loop and
|
||||
* releases the memory and verifies the meminfo
|
||||
*
|
||||
* In both cases, the memory info before allocation and
|
||||
* after releasing the memory should be the same
|
||||
*
|
||||
*/
|
||||
template<typename T>
|
||||
static void MemoryAllocDiffSizes(int gpu) {
|
||||
HIP_CHECK(hipSetDevice(gpu));
|
||||
std::vector<size_t> array_size;
|
||||
array_size.push_back(SMALLCHUNK_NUMH);
|
||||
array_size.push_back(LARGECHUNK_NUMH);
|
||||
for (auto &sizes : array_size) {
|
||||
T* A_d[CHUNK_LOOP];
|
||||
size_t pitch_A;
|
||||
size_t width;
|
||||
if (sizes == SMALLCHUNK_NUMH) {
|
||||
width = SMALLCHUNK_NUMW * sizeof(T);
|
||||
} else {
|
||||
width = LARGECHUNK_NUMW * sizeof(T);
|
||||
}
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < CHUNK_LOOP; i++) {
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d[i]),
|
||||
&pitch_A, width, sizes));
|
||||
}
|
||||
for (int i = 0; i < CHUNK_LOOP; i++) {
|
||||
HIP_CHECK(hipFree(A_d[i]));
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
if (pavail != avail) {
|
||||
HIPASSERT(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/*Thread Function */
|
||||
static void threadFunc(int gpu) {
|
||||
MemoryAllocDiffSizes<float>(gpu);
|
||||
}
|
||||
/*
|
||||
* This testcase verifies the negative scenarios of hipMallocPitch API
|
||||
*/
|
||||
TEST_CASE("Unit_hipMallocPitch_Negative") {
|
||||
float* A_d;
|
||||
size_t pitch_A;
|
||||
size_t width{NUM_W * sizeof(float)};
|
||||
#if HT_NVIDIA
|
||||
SECTION("NullPtr to Pitched Ptr") {
|
||||
REQUIRE(hipMallocPitch(nullptr,
|
||||
&pitch_A, width, NUM_H) != hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("nullptr to pitch") {
|
||||
REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
nullptr, width, NUM_H) != hipSuccess);
|
||||
}
|
||||
#endif
|
||||
SECTION("Width 0 in hipMallocPitch") {
|
||||
REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, 0, NUM_H) == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Height 0 in hipMallocPitch") {
|
||||
REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, 0) == hipSuccess);
|
||||
}
|
||||
|
||||
SECTION("Max int values") {
|
||||
REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, std::numeric_limits<int>::max(),
|
||||
std::numeric_limits<int>::max()) != hipSuccess);
|
||||
}
|
||||
}
|
||||
/*
|
||||
* This testcase verifies the basic scenario of
|
||||
* hipMallocPitch API for different datatypes
|
||||
*
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic",
|
||||
"[hipMallocPitch]", int, unsigned int, float) {
|
||||
TestType* A_d;
|
||||
size_t pitch_A;
|
||||
size_t width{NUM_W * sizeof(TestType)};
|
||||
REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, NUM_H) == hipSuccess);
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
}
|
||||
|
||||
/*
|
||||
* This testcase verifies hipMallocPitch API for small
|
||||
* and big chunks of data.
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipMallocPitch_SmallandBigChunks",
|
||||
"[hipMallocPitch]", int, unsigned int, float) {
|
||||
MemoryAllocDiffSizes<TestType>(0);
|
||||
}
|
||||
|
||||
/*
|
||||
* This testcase verifies the memory allocated by hipMallocPitch API
|
||||
* by performing Memcpy2D on the allocated memory.
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Memcpy2D", ""
|
||||
, int, float, double) {
|
||||
HIP_CHECK(hipSetDevice(0));
|
||||
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
|
||||
*B_d{nullptr};
|
||||
size_t pitch_A, pitch_B;
|
||||
size_t width{NUM_W * sizeof(TestType)};
|
||||
|
||||
// Allocating memory
|
||||
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
|
||||
&A_h, &B_h, &C_h, NUM_W*NUM_H, false);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, NUM_H));
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
|
||||
&pitch_B, width, NUM_H));
|
||||
|
||||
// Initialize the data
|
||||
HipTest::setDefaultData<TestType>(NUM_W*NUM_H, A_h, B_h, C_h);
|
||||
|
||||
// Host to Device
|
||||
HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
|
||||
COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice));
|
||||
|
||||
// Performs D2D on same GPU device
|
||||
HIP_CHECK(hipMemcpy2D(B_d, pitch_B, A_d,
|
||||
pitch_A, COLUMNS*sizeof(TestType),
|
||||
ROWS, hipMemcpyDeviceToDevice));
|
||||
|
||||
// hipMemcpy2D Device to Host
|
||||
HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
|
||||
COLUMNS*sizeof(TestType), ROWS,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
// Validating the result
|
||||
REQUIRE(HipTest::checkArray<TestType>(A_h, B_h, COLUMNS, ROWS) == true);
|
||||
|
||||
|
||||
// DeAllocating the memory
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(B_d));
|
||||
HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
|
||||
A_h, B_h, C_h, false);
|
||||
}
|
||||
|
||||
|
||||
|
||||
/*
|
||||
This testcase verifies the hipMallocPitch API in multithreaded
|
||||
scenario by launching threads in parallel on multiple GPUs
|
||||
and verifies the hipMallocPitch API with small and big chunks data
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_hipMallocPitch_MultiThread", "") {
|
||||
std::vector<std::thread> threadlist;
|
||||
int devCnt = 0;
|
||||
|
||||
devCnt = HipTest::getDeviceCount();
|
||||
|
||||
size_t tot, avail, ptot, pavail;
|
||||
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
|
||||
for (int i = 0; i < devCnt; i++) {
|
||||
threadlist.push_back(std::thread(threadFunc, i));
|
||||
}
|
||||
|
||||
for (auto &t : threadlist) {
|
||||
t.join();
|
||||
}
|
||||
HIP_CHECK(hipMemGetInfo(&avail, &tot));
|
||||
|
||||
if (pavail != avail) {
|
||||
WARN("Memory leak of hipMallocPitch API in multithreaded scenario");
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
/*
|
||||
* This testcase verifies hipMallocPitch API by
|
||||
* 1. Allocating Memory using hipMallocPitch API
|
||||
* 2. Launching the kernel and copying the data from the allocated kernel
|
||||
* variable to another kernel variable.
|
||||
* 3. Validating the result
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", ""
|
||||
, int, float, double) {
|
||||
HIP_CHECK(hipSetDevice(0));
|
||||
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr},
|
||||
*B_d{nullptr};
|
||||
size_t pitch_A, pitch_B;
|
||||
size_t width{NUM_W * sizeof(TestType)};
|
||||
|
||||
// Allocating memory
|
||||
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
|
||||
&A_h, &B_h, &C_h, NUM_W*NUM_H, false);
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d),
|
||||
&pitch_A, width, NUM_H));
|
||||
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&B_d),
|
||||
&pitch_B, width, NUM_H));
|
||||
|
||||
// Host to Device
|
||||
HIP_CHECK(hipMemcpy2D(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType),
|
||||
COLUMNS*sizeof(TestType), ROWS, hipMemcpyHostToDevice));
|
||||
|
||||
|
||||
hipLaunchKernelGGL(copy_var<TestType>, dim3(1), dim3(1),
|
||||
0, 0, static_cast<TestType*>(A_d),
|
||||
static_cast<TestType*>(B_d), ROWS, pitch_A);
|
||||
|
||||
|
||||
// hipMemcpy2D Device to Host
|
||||
HIP_CHECK(hipMemcpy2D(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B,
|
||||
COLUMNS*sizeof(TestType), ROWS,
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
// Validating the result
|
||||
validateResult(A_h, B_h, pitch_A);
|
||||
|
||||
// DeAllocating the memory
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(B_d));
|
||||
HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
|
||||
A_h, B_h, C_h, false);
|
||||
}
|
||||
|
||||
|
||||
Ссылка в новой задаче
Block a user