AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: Icc24ccd7e75589d0665c7eea87ccafcec17634d7


[ROCm/hip commit: 483dd2595d]
This commit is contained in:
Jenkins
2022-05-25 19:10:38 -04:00
parent 3fd6792649 5298b0b4c4
commit dafa216c3c
14 changed files with 938 additions and 652 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/catch/hipTestMain/config/config_amd_linux.json
+2
View File
@@ -1,6 +1,8 @@
{
"DisabledTests":
[
"# Following test is related to ticket EXSWCPHIPT-41",
"Unit_hipStreamGetPriority_happy",
"Unit_hipStreamPerThread_DeviceReset_1"
]
projects/hip/tests/catch/hipTestMain/config/config_amd_windows.json
+3 -2
View File
@@ -73,7 +73,8 @@
"Unit_hipGraphAddEventRecordNode_Functional_WithFlags",
"Unit_hipGraphAddEventRecordNode_MultipleRun",
"Unit_hipMalloc3D_Negative",
"Unit_hipPointerGetAttribute_MappedMem"
"Unit_hipPointerGetAttribute_MappedMem",
"# Following test is related to ticket EXSWCPHIPT-41",
"Unit_hipStreamGetPriority_happy"
]
}
projects/hip/tests/catch/include/hip_test_common.hh
+38 -13
View File
@@ -27,6 +27,7 @@ THE SOFTWARE.
#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
// Not thread-safe
#define HIP_CHECK(error) \
{ \
hipError_t localError = error; \
@@ -37,6 +38,20 @@ THE SOFTWARE.
} \
}
// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError.
#define HIP_CHECK_ERROR(errorExpr, expectedError) \
{ \
hipError_t localError = errorExpr; \
INFO("Matching Errors: " \
<< " Expected Error: " << hipGetErrorString(expectedError) \
<< " Expected Code: " << expectedError << '\n' \
<< " Actual Error: " << hipGetErrorString(localError) \
<< " Actual Code: " << localError << "\nStr: " << #errorExpr \
<< "\nIn File: " << __FILE__ << " At line: " << __LINE__); \
REQUIRE(localError == expectedError); \
}
// Not thread-safe
#define HIPRTC_CHECK(error) \
{ \
auto localError = error; \
@@ -46,25 +61,26 @@ THE SOFTWARE.
REQUIRE(false); \
} \
}
// Although its assert, it will be evaluated at runtime
#define HIP_ASSERT(x) \
{ REQUIRE((x)); }
#ifdef __cplusplus
#include <iostream>
#include <iomanip>
#include <chrono>
#include <iostream>
#include <iomanip>
#include <chrono>
#endif
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), \
localError, #error, __FILE__, __LINE__); \
abort(); \
} \
}
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), localError, \
#error, __FILE__, __LINE__); \
abort(); \
} \
}
#define HIPASSERT(condition) \
if (!(condition)) { \
@@ -104,8 +120,8 @@ static inline int getDeviceCount() {
// Returns the current system time in microseconds
static inline long long get_time() {
return std::chrono::high_resolution_clock::now().time_since_epoch()
/std::chrono::microseconds(1);
return std::chrono::high_resolution_clock::now().time_since_epoch() /
std::chrono::microseconds(1);
}
static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) {
@@ -147,7 +163,16 @@ inline bool isImageSupported() {
return imageSupport != 0;
}
/**
* Causes the test to stop and be skipped at runtime.
* reason: Message describing the reason the test has been skipped.
*/
static inline void HIP_SKIP_TEST(char const* const reason) noexcept {
// ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test.
std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl;
}
}
// This must be called in the beginning of image test app's main() to indicate whether image
// is supported.
projects/hip/tests/catch/stress/memory/hipMalloc.cc
+50
View File
@@ -0,0 +1,50 @@
/*
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 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.
*/
#include <hip_test_common.hh>
#include <memory>
// Stress allocation tests
// Try to allocate as much memory as possible
// But since max allocation can fail, we need to be happy with atleast 1/4th of memory
TEST_CASE("Stress_hipMalloc_HighSizeAlloc") {
size_t devMemTotal{0}, devMemFree{0};
HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemTotal));
REQUIRE(devMemFree > 0);
REQUIRE(devMemTotal > 0);
char* d_ptr{nullptr};
size_t counter{0};
INFO("Free Mem Available: " << devMemFree << " bytes out of " << devMemTotal << " bytes!");
while (hipMalloc(&d_ptr, devMemFree) != hipSuccess && devMemFree > 1) {
counter++;
devMemFree >>= 1; // reduce the memory to be allocated by half
INFO("Attempt to allocate " << devMemFree << " bytes out of " << devMemTotal
<< " bytes failed!");
REQUIRE(counter <= 2); // Make sure that we are atleast able to allocate 1/4th of max memory
}
HIP_CHECK(hipMemset(d_ptr, 1, devMemFree));
auto ptr = std::unique_ptr<unsigned char[]>{new unsigned char[devMemFree]};
HIP_CHECK(hipMemcpy(ptr.get(), d_ptr, devMemFree, hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(d_ptr));
REQUIRE(std::all_of(ptr.get(), ptr.get() + devMemFree, [](unsigned char n) { return n == 1; }));
}
projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc
+4 -4
View File
@@ -150,7 +150,7 @@ static int HmmAttrPrint() {
// The following test case allocation, host access, device access of HMM
// memory from size 1 to 10KB
TEST_CASE("Unit_hipMallocManaged_MultiSize") {
TEST_CASE("Stress_hipMallocManaged_MultiSize") {
IfTestPassed = true;
int managed = HmmAttrPrint();
if (managed == 1) {
@@ -196,7 +196,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") {
// The following test case tests the behavior of kernel with a HMM memory and
// hipMalloc memory
TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
TEST_CASE("Stress_hipMallocManaged_KrnlWth2MemTypes") {
IfTestPassed = true;
int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
size_t NumElms = (1024 * 1024);
@@ -241,7 +241,7 @@ TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
// The following test case tests when the same Hmm memory is used for
// launching multiple different kernels will results in any issue
TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
TEST_CASE("Stress_hipMallocManaged_MultiKrnlHmmAccess") {
int managed = HmmAttrPrint();
if (managed) {
int InitVal = 123, NumElms = (1024 * 1024);
@@ -253,7 +253,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
}
// Testing the allocation of/scenarios around max possible memory
TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
TEST_CASE("Stress_hipMallocManaged_ExtremeSizes") {
int managed = HmmAttrPrint();
if (managed == 1) {
bool IfTestPassed = true;
projects/hip/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc
+165 -126
View File
@@ -26,117 +26,131 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define NUM 1024
#define SIZE 1024 * 4
constexpr size_t NUM = 1024;
constexpr size_t SIZE = 1024 * 4;
__device__ int globalIn[NUM];
__device__ int globalOut[NUM];
__global__ void Assign(int* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Out[tid] = globalIn[tid];
globalOut[tid] = globalIn[tid];
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Out[tid] = globalIn[tid];
globalOut[tid] = globalIn[tid];
}
__device__ __constant__ int globalConstArr[NUM];
__device__ __constant__ static float statConstVar = 1.0f;
__device__ __constant__ int globalConst[NUM];
__global__ void checkGlobalConstAddress(int* addr, bool* out) {
*out = (globalConstArr == addr);
}
__global__ void checkAddress(int* addr, bool* out) { *out = (globalConst == addr); }
__global__ void checkStaticConstVarAddress(float* addr, bool* out) {
*out = (&statConstVar == addr);
}
/**
Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread.
*/
TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
int *A, *Am, *B, *Ad, *C, *Cm;
A = new int[NUM];
B = new int[NUM];
C = new int[NUM];
for (int i = 0; i < NUM; i++) {
A[i] = -1 * i;
B[i] = 0;
C[i] = 0;
}
int *A{nullptr}, *Am{nullptr}, *B{nullptr}, *Ad{nullptr}, *C{nullptr}, *Cm{nullptr};
A = new int[NUM];
B = new int[NUM];
C = new int[NUM];
HIP_CHECK(hipMalloc(&Ad, SIZE));
HIP_CHECK(hipHostMalloc(&Am, SIZE));
HIP_CHECK(hipHostMalloc(&Cm, SIZE));
for (int i = 0; i < NUM; i++) {
Am[i] = -1 * i;
Cm[i] = 0;
}
HIP_CHECK(hipMalloc((void**)&Ad, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Am, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE));
hipStream_t stream;
for (size_t i = 0; i < NUM; i++) {
A[i] = -1 * static_cast<int>(i);
B[i] = 0;
C[i] = 0;
Am[i] = -1 * static_cast<int>(i);
Cm[i] = 0;
}
SECTION("Calling hipMemcpyTo/FromSymbol using stream") {
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (int i = 0; i < NUM; i++) {
REQUIRE(Am[i] == B[i]);
REQUIRE(Am[i] == Cm[i]);
}
for (int i = 0; i < NUM; i++) {
A[i] = -2 * i;
B[i] = 0;
}
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost));
for (int i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
for (int i = 0; i < NUM; i++) {
A[i] = -3 * i;
B[i] = 0;
}
SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
for (int i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipHostFree(Am));
HIP_CHECK(hipHostFree(Cm));
HIP_CHECK(hipFree(Ad));
delete[] A;
delete[] B;
delete[] C;
for (size_t i = 0; i < NUM; i++) {
REQUIRE(Am[i] == B[i]);
REQUIRE(Am[i] == Cm[i]);
}
}
SECTION("Calling hipMemcpyTo/FromSymbol - validate value in host memory") {
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost));
for (size_t i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
}
SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") {
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(
hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
for (size_t i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
}
SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice,
hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
for (size_t i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
}
// Check for address on GPU and CPU side and compare it
// If address mismatch report error
// Validate size of symbol as well, compare it with output of hipGetSymbolSize
SECTION("Validate address on GPU") {
bool* checkOkD{nullptr};
bool checkOk = false;
size_t symbolSize = 0;
int* symbolAddress{nullptr};
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipGetSymbolAddress((void**)&symbolAddress, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT((symbolSize == SIZE));
}
HIP_CHECK(hipHostFree(Am));
HIP_CHECK(hipHostFree(Cm));
HIP_CHECK(hipFree(Ad));
delete[] A;
delete[] B;
delete[] C;
}
/**
@@ -144,39 +158,64 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
2) Validate get symbol address/size for static const variable.
*/
TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") {
bool *checkOkD;
bool checkOk = false;
size_t symbolSize{};
int *symbolArrAddress{};
float *symbolVarAddress{};
bool* checkOkD{nullptr};
bool checkOk = false;
size_t symbolSize{};
int* symbolArrAddress{};
float* symbolVarAddress{};
SECTION("Validate symbol size/address of global const array") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipGetSymbolAddress(
reinterpret_cast<void **>(&symbolArrAddress),
HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1),
0, 0, symbolArrAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == SIZE);
}
SECTION("Validate symbol size/address of global const array") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipGetSymbolAddress(reinterpret_cast<void**>(&symbolArrAddress),
HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
symbolArrAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == SIZE);
}
SECTION("Validate symbol size/address of static const variable") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipGetSymbolAddress(
reinterpret_cast<void **>(&symbolVarAddress),
HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1),
dim3(1, 1, 1), 0, 0, symbolVarAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == sizeof(float));
}
SECTION("Validate symbol size/address of static const variable") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
HIP_CHECK(
hipGetSymbolAddress(reinterpret_cast<void**>(&symbolVarAddress), HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
symbolVarAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == sizeof(float));
}
}
TEST_CASE("Unit_hipGetSymbolAddress_Negative") {
SECTION("Invalid symbol") {
int notADeviceSymbol{0};
int* addr{nullptr};
HIP_CHECK_ERROR(
hipGetSymbolAddress(reinterpret_cast<void**>(&addr), HIP_SYMBOL(notADeviceSymbol)),
hipErrorInvalidSymbol);
}
SECTION("Nullptr symbol") {
int* addr{nullptr};
HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&addr), nullptr),
hipErrorInvalidSymbol);
}
}
TEST_CASE("Unit_hipGetSymbolSize_Negative") {
SECTION("Invalid symbol") {
int notADeviceSymbol{0};
size_t dsize{0};
HIP_CHECK_ERROR(hipGetSymbolSize(&dsize, HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol);
}
SECTION("Nullptr symbol") {
size_t size{0};
HIP_CHECK_ERROR(hipGetSymbolSize(&size, nullptr), hipErrorInvalidSymbol);
}
}
projects/hip/tests/catch/unit/memory/CMakeLists.txt
+2 -4
View File
@@ -42,14 +42,13 @@ set(TEST_SRC
hipPointerGetAttributes.cc
hipHostGetFlags.cc
hipMallocManaged_MultiScenario.cc
hipMemsetInvalidPtr.cc
hipMemsetNegative.cc
hipMemset.cc
hipMemsetAsyncMultiThread.cc
hipMemset3D.cc
hipMemset2D.cc
hipHostMallocTests.cc
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
hipMallocManagedFlagsTst.cc
hipMemPrefetchAsyncExtTsts.cc
@@ -108,14 +107,13 @@ set(TEST_SRC
hipHostRegister.cc
hipHostGetFlags.cc
hipMallocManaged_MultiScenario.cc
hipMemsetInvalidPtr.cc
hipMemsetNegative.cc
hipMemset.cc
hipMemsetAsyncMultiThread.cc
hipMemset3D.cc
hipMemset2D.cc
hipHostMallocTests.cc
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
hipMallocManagedFlagsTst.cc
hipMemPrefetchAsyncExtTsts.cc
projects/hip/tests/catch/unit/memory/hipMallocArray.cc
+276 -47
View File
@@ -25,7 +25,6 @@ hipMallocArray API test scenarios
4. Multithreaded scenario
*/
#include <hip_test_common.hh>
static constexpr auto NUM_W{4};
@@ -53,20 +52,16 @@ static void MallocArray_DiffSizes(int gpu) {
std::vector<size_t> array_size;
array_size.push_back(NUM_W);
array_size.push_back(BIGNUM_W);
for (auto &size : array_size) {
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));
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));
HIP_CHECK(hipMallocArray(&A_d[i], &desc, BIGNUM_W, BIGNUM_H, hipArrayDefault));
}
}
for (int i = 0; i < ARRAY_LOOP; i++) {
@@ -79,10 +74,6 @@ static void MallocArray_DiffSizes(int gpu) {
}
}
static void MallocArrayThreadFunc(int gpu) {
MallocArray_DiffSizes(gpu);
}
/*
* This testcase verifies the negative scenarios of
* hipMallocArray API
@@ -92,56 +83,31 @@ TEST_CASE("Unit_hipMallocArray_Negative") {
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
#if HT_NVIDIA
SECTION("NullPointer to Array") {
REQUIRE(hipMallocArray(nullptr, &desc,
NUM_W, NUM_H, hipArrayDefault) != hipSuccess);
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);
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);
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);
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("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);
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);
}
TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); }
/*
@@ -156,10 +122,10 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
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));
threadlist.push_back(std::thread(MallocArray_DiffSizes, i));
}
for (auto &t : threadlist) {
for (auto& t : threadlist) {
t.join();
}
HIP_CHECK(hipMemGetInfo(&avail, &tot));
@@ -170,3 +136,266 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") {
}
}
constexpr size_t BlockSize = 16;
template <class T, size_t N> struct type_and_size {
using type = T;
static constexpr size_t size = N;
};
// scalars are interpreted as a vector of 1 length.
// template <size_t N> using int_constant = std::integral_constant<size_t, N>;
template <typename T> struct vector_info;
template <> struct vector_info<int> : type_and_size<int, 1> {};
template <> struct vector_info<float> : type_and_size<float, 1> {};
template <> struct vector_info<short> : type_and_size<short, 1> {};
template <> struct vector_info<char> : type_and_size<char, 1> {};
template <> struct vector_info<unsigned int> : type_and_size<unsigned int, 1> {};
template <> struct vector_info<unsigned short> : type_and_size<unsigned short, 1> {};
template <> struct vector_info<unsigned char> : type_and_size<unsigned char, 1> {};
template <> struct vector_info<int2> : type_and_size<int, 2> {};
template <> struct vector_info<float2> : type_and_size<float, 2> {};
template <> struct vector_info<short2> : type_and_size<short, 2> {};
template <> struct vector_info<char2> : type_and_size<char, 2> {};
template <> struct vector_info<uint2> : type_and_size<unsigned int, 2> {};
template <> struct vector_info<ushort2> : type_and_size<unsigned short, 2> {};
template <> struct vector_info<uchar2> : type_and_size<unsigned char, 2> {};
template <> struct vector_info<int4> : type_and_size<int, 4> {};
template <> struct vector_info<float4> : type_and_size<float, 4> {};
template <> struct vector_info<short4> : type_and_size<short, 4> {};
template <> struct vector_info<char4> : type_and_size<char, 4> {};
template <> struct vector_info<uint4> : type_and_size<unsigned int, 4> {};
template <> struct vector_info<ushort4> : type_and_size<unsigned short, 4> {};
template <> struct vector_info<uchar4> : type_and_size<unsigned char, 4> {};
// Kernels ///////////////////////////////////////
// read from a texture using normalized coordinates
constexpr size_t ChannelToRead = 1;
template <typename T>
__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
bool textureGather) {
// Calculate normalized texture coordinates
const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
const float u = x / (float)width;
// Read from texture and write to global memory
if (height == 0) {
output[x] = tex1D<T>(texObj, u);
} else {
const float v = y / (float)height;
output[y * width + x] =
textureGather ? tex2Dgather<T>(texObj, u, v, ChannelToRead) : tex2D<T>(texObj, u, v);
}
}
template <typename T> __device__ void addOne(T* a) {
using scalar_type = typename vector_info<T>::type;
auto as = reinterpret_cast<scalar_type*>(a);
for (size_t i = 0; i < vector_info<T>::size; ++i) {
as[i] = as[i] + static_cast<scalar_type>(1);
}
}
// read from a surface and write to another
template <typename T> __global__ void incSurface(hipSurfaceObject_t surf, size_t height) {
// Calculate surface coordinates
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if (height == 0) {
T data;
surf1Dread(&data, surf, x * sizeof(T));
addOne(&data); // change the value to show that write works
surf1Dwrite(data, surf, x * sizeof(T));
} else {
T data;
surf2Dread(&data, surf, x * sizeof(T), y);
addOne(&data); // change the value to show that write works
surf2Dwrite(data, surf, x * sizeof(T), y);
}
}
// Helpers ///////////////////////////////////////
template <typename T> size_t getAllocSize(const size_t width, const size_t height) noexcept {
return sizeof(T) * width * (height ? height : 1);
}
template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {
bool allMatch = true;
size_t i = 0;
for (; i < hostData.size(); ++i) {
allMatch = allMatch && hostData[i] == static_cast<T>(i);
if (!allMatch) break;
}
INFO("hostData[" << i << "] == " << static_cast<T>(hostData[i]));
REQUIRE(allMatch);
}
// Tests /////////////////////////////////////////
// Test the default array by generating a texture from it then reading from that texture.
// Textures are read-only so write to the array then copy from the texture into normal device memory
template <typename T>
void testArrayAsTexture(hipArray_t arrayPtr, const size_t width, const size_t height) {
using scalar_type = typename vector_info<T>::type;
constexpr auto vec_size = vector_info<T>::size;
const auto h = height ? height : 1;
const size_t pitch = width * sizeof(T); // no padding
const auto size = pitch * h;
// create an array to initialize the hip array, then later use it to hold the result
std::vector<scalar_type> hostData(width * h * vec_size);
// Setup backing array
// assign ascending values to the data array to show indexing is working.
std::iota(std::begin(hostData), std::end(hostData), 0);
HIP_CHECK(
hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice));
// create texture
hipTextureObject_t textObj{};
hipResourceDesc resDesc{};
memset(&resDesc, 0, sizeof(hipResourceDesc));
// enum to store how to resDesc.res union is being used
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = arrayPtr;
hipTextureDesc textDesc{};
memset(&textDesc, 0, sizeof(hipTextureDesc));
textDesc.filterMode =
hipFilterModePoint; // use the actual values in the texture, not normalized data
textDesc.readMode = hipReadModeElementType; // don't convert the data to floats
textDesc.normalizedCoords = 1; // use normalized coordinates (0.0-1.0)
HIP_CHECK(hipCreateTextureObject(&textObj, &resDesc, &textDesc, nullptr));
// run kernel
T* device_data{};
HIP_CHECK(hipMalloc(&device_data, size));
readFromTexture<<<dim3(width / BlockSize, height ? height / BlockSize : 1, 1),
dim3(BlockSize, height ? BlockSize : 1, 1)>>>(device_data, textObj, width,
height, false);
HIP_CHECK(hipGetLastError()); // check for errors when running the kernel
// copy data back and then test it
std::fill(std::begin(hostData), std::end(hostData), 0);
HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost));
checkDataIsAscending(hostData);
// clean up
HIP_CHECK(hipDestroyTextureObject(textObj));
HIP_CHECK(hipFree(device_data));
}
// Test the an array created with the SurfaceLoadStore flag by generating a surface and reading from
// it and writing to it.
template <typename T>
void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t height) {
using scalar_type = typename vector_info<T>::type;
constexpr auto vec_size = vector_info<T>::size;
const auto h = height ? height : 1;
const size_t pitch = width * sizeof(T); // no padding
const auto size = pitch * h;
std::vector<scalar_type> hostData(width * h * vec_size);
// Setup backing array
// assign ascending values to the data array to show indexing is working.
std::iota(std::begin(hostData), std::end(hostData), 0);
HIP_CHECK(
hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice));
// create surface
hipSurfaceObject_t surfObj{};
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(hipResourceDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = arrayPtr;
HIP_CHECK(hipCreateSurfaceObject(&surfObj, &resDesc));
// run kernel
T* device_data{};
HIP_CHECK(hipMalloc(&device_data, size));
// This will increment the values of the surface, so this is undone later
incSurface<T><<<dim3(width / BlockSize, height ? height / BlockSize : 1, 1),
dim3(BlockSize, height ? BlockSize : 1, 1)>>>(surfObj, height);
HIP_CHECK(hipGetLastError()); // check for errors when running the kernel
// copy data back and then test it
std::fill(std::begin(hostData), std::end(hostData), 0);
HIP_CHECK(hipMemcpy2DFromArray(hostData.data(), pitch, arrayPtr, 0, 0, pitch, h,
hipMemcpyDeviceToHost));
// undo the increment
std::for_each(std::begin(hostData), std::end(hostData),
[](scalar_type& x) { x -= static_cast<scalar_type>(1); });
checkDataIsAscending(hostData);
// clean up
HIP_CHECK(hipDestroySurfaceObject(surfObj));
HIP_CHECK(hipFree(device_data));
}
size_t getFreeMem() {
size_t free = 0, total = 0;
HIP_CHECK(hipMemGetInfo(&free, &total));
return free;
}
// The happy path of a default array and a SurfaceLoadStore array should work
// Selection of types chosen to reduce compile times
TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, short2, char, uchar2,
char4, float, float2, float4) {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-62");
#endif
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
size_t init_free = getFreeMem();
// pointer to the array in device memory
hipArray_t arrayPtr{};
size_t width = 1024;
size_t height = GENERATE(0, 1024);
SECTION("hipArrayDefault") {
INFO("flag is hipArrayDefault");
INFO("height: " << height);
HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault));
testArrayAsTexture<TestType>(arrayPtr, width, height);
}
#if HT_NVIDIA // surfaces and texture gather not supported on AMD
SECTION("hipArraySurfaceLoadStore") {
INFO("flag is hipArraySurfaceLoadStore");
INFO("height: " << height);
HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArraySurfaceLoadStore));
testArrayAsSurface<TestType>(arrayPtr, width, height);
}
#endif
size_t final_free = getFreeMem();
const size_t alloc_size = getAllocSize<TestType>(width, height);
// alloc will be chunked, so this is not exact
REQUIRE(init_free - final_free >= alloc_size);
HIP_CHECK(hipFreeArray(arrayPtr));
}
projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc
+48 -59
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -51,14 +51,13 @@ Testcase Scenarios :
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include <vector>
#include <limits>
#include <atomic>
#include <limits>
#include <vector>
/* Buffer size for bigger chunks in alloc/free cycles */
static constexpr auto BuffSizeBC = 5*1024*1024;
static constexpr auto BuffSizeBC = 5 * 1024 * 1024;
/* Buffer size for smaller chunks in alloc/free cycles */
static constexpr auto BuffSizeSC = 16;
@@ -68,19 +67,18 @@ static constexpr auto BuffSizeSC = 16;
static constexpr auto NumDiv = 100;
/* Max alloc/free iterations for smaller chunks */
static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv);
static constexpr auto MaxAllocFree_SmallChunks = (5000000 / NumDiv);
/* Max alloc/free iterations for bigger chunks */
static constexpr auto MaxAllocFree_BigChunks = 10000;
/* Max alloc and pool iterations */
static constexpr auto MaxAllocPoolIter = (2000000/NumDiv);
static constexpr auto MaxAllocPoolIter = (2000000 / NumDiv);
/* Test status shared across threads */
static std::atomic<bool> g_thTestPassed{true};
/**
* Validates data consistency on supplied gpu
*/
@@ -103,9 +101,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
@@ -121,9 +118,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
// In concurrent calls on one GPU, we cannot verify leaking in this way
UNSCOPED_INFO(
"validateMemoryOnGPU : Memory allocation mismatch observed."
<< "Possible memory leak.");
UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed."
<< "Possible memory leak.");
TestPassed = false;
}
@@ -138,7 +134,7 @@ static bool regressAllocInLoop(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
int i = 0;
int *ptr;
int* ptr;
HIP_CHECK(hipSetDevice(gpu));
numBytes = BuffSizeBC;
@@ -150,11 +146,12 @@ static bool regressAllocInLoop(int gpu) {
HIP_CHECK(hipMemGetInfo(&avail, &tot));
HIP_CHECK(hipFree(ptr));
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
tot << ", pavail-avail=" << pavail-avail);
if (pavail - avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes
<< " not matching with hipMemGetInfo - FAIL."
<< "pavail=" << pavail << ", ptot=" << ptot
<< ", avail=" << avail << ", tot=" << tot
<< ", pavail-avail=" << pavail - avail);
TestPassed = false;
break;
}
@@ -173,8 +170,8 @@ static bool regressAllocInLoop(int gpu) {
HIP_CHECK(hipMemGetInfo(&avail, &tot));
if ((pavail != avail) || (ptot != tot)) {
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
"Possible memory leak.");
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed."
<< "Possible memory leak.");
TestPassed &= false;
}
@@ -203,9 +200,8 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
@@ -238,7 +234,7 @@ static bool regressAllocInLoopMthread(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
int i = 0;
int *ptr;
int* ptr;
HIPCHECK(hipSetDevice(gpu));
numBytes = BuffSizeBC;
@@ -250,11 +246,12 @@ static bool regressAllocInLoopMthread(int gpu) {
HIPCHECK(hipMemGetInfo(&avail, &tot));
HIPCHECK(hipFree(ptr));
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
tot << ", pavail-avail=" << pavail-avail);
if (pavail - avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes
<< " not matching with hipMemGetInfo - FAIL."
<< "pavail=" << pavail << ", ptot=" << ptot
<< ", avail=" << avail << ", tot=" << tot
<< ", pavail-avail=" << pavail - avail);
TestPassed = false;
break;
}
@@ -273,8 +270,8 @@ static bool regressAllocInLoopMthread(int gpu) {
HIPCHECK(hipMemGetInfo(&avail, &tot));
if ((pavail != avail) || (ptot != tot)) {
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
"Possible memory leak.");
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed."
<< "Possible memory leak.");
TestPassed &= false;
}
@@ -285,18 +282,15 @@ static bool regressAllocInLoopMthread(int gpu) {
* Thread func to regress alloc and check data consistency
*/
static void threadFunc(int gpu) {
g_thTestPassed = regressAllocInLoopMthread(gpu)
&& validateMemoryOnGpuMThread(gpu);
g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu);
UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" <<
g_thTestPassed.load());
UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load());
}
/* Performs Argument Validation of api */
TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
int *ptr;
hipError_t ret;
int* ptr{nullptr};
SECTION("hipMalloc() when size(0)") {
HIP_CHECK(hipMalloc(&ptr, 0));
@@ -304,21 +298,17 @@ TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
REQUIRE(ptr == nullptr);
}
SECTION("hipFree() when freeing nullptr ") {
ptr = nullptr;
// api should return success and shudnt crash
SECTION("hipFree() when freeing nullptr") {
HIP_CHECK(hipFree(ptr));
}
SECTION("hipMalloc() with invalid argument") {
constexpr auto sizeBytes = 100;
ret = hipMalloc(nullptr, sizeBytes);
REQUIRE(ret != hipSuccess);
HIP_CHECK_ERROR(hipMalloc(nullptr, 100), hipErrorInvalidValue);
}
SECTION("hipMalloc() with max size_t") {
ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max());
REQUIRE(ret != hipSuccess);
HIP_CHECK_ERROR(hipMalloc(&ptr, std::numeric_limits<std::size_t>::max()),
hipErrorMemoryAllocation);
}
}
@@ -344,12 +334,12 @@ TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
* of time.
*/
TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
size_t avail, tot, pavail, ptot;
bool ret;
hipError_t err;
std::vector<int *> ptrlist;
size_t avail{0}, tot{0}, pavail{0}, ptot{0};
bool ret{false};
hipError_t err{};
std::vector<int*> ptrlist{};
constexpr auto BuffSize = 10;
int devCnt, *ptr;
int devCnt{0}, *ptr{nullptr};
// Get GPU count
HIP_CHECK(hipGetDeviceCount(&devCnt));
@@ -358,14 +348,13 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
// Allocate small chunks of memory million times
for (int i = 0; i < MaxAllocPoolIter ; i++) {
for (int i = 0; i < MaxAllocPoolIter; i++) {
if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) {
HIP_CHECK(hipMemGetInfo(&avail, &tot));
INFO("Loop regression pool allocation failure. " <<
"Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " <<
avail/(1024.0*1024.0) << " iter " << i << " error "
<< hipGetErrorString(err));
INFO("Loop regression pool allocation failure. "
<< "Total gpu memory " << tot / (1024.0 * 1024.0) << ", Free memory "
<< avail / (1024.0 * 1024.0) << " iter " << i << " error " << hipGetErrorString(err));
REQUIRE(false);
}
@@ -375,7 +364,7 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
}
// Free ptrs at later point of time
for ( auto &t : ptrlist ) {
for (auto& t : ptrlist) {
HIP_CHECK(hipFree(t));
}
@@ -404,7 +393,7 @@ TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") {
threadlist.push_back(std::thread(threadFunc, i));
}
for (auto &t : threadlist) {
for (auto& t : threadlist) {
t.join();
}
projects/hip/tests/catch/unit/memory/hipMemset3DNegative.cc
-236
View File
@@ -1,236 +0,0 @@
/*
Copyright (c) 2021 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 Scenarios :
1) Test hipMemset3D() with uninitialized devPitchedPtr.
2) Test hipMemset3DAsync() with uninitialized devPitchedPtr.
3) Reset devPitchedPtr to zero and check return value for hipMemset3D().
4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync().
5) Test hipMemset3D() with extent.width as max size_t and keeping height,
depth as valid values.
6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height,
depth as valid values.
7) Test hipMemset3D() with extent.height as max size_t and keeping width,
depth as valid values.
8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width,
depth as valid values.
9) Test hipMemset3D() with extent.depth as max size_t and keeping height,
width as valid values.
10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height,
width as valid values.
11) Device Ptr out bound and extent(0) passed for hipMemset3D().
12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync().
13) Device Ptr out bound and valid extent passed for hipMemset3D().
14) Device Ptr out bound and valid extent passed for hipMemset3DAsync().
*/
#include <hip_test_common.hh>
TEST_CASE("Unit_hipMemset3D_Negative") {
hipError_t ret;
hipPitchedPtr devPitchedPtr;
constexpr int memsetval = 1;
constexpr size_t numH = 256, numW = 256;
constexpr size_t depth = 10;
constexpr size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
SECTION("Using uninitialized devpitched ptr") {
hipPitchedPtr devPitchedUnPtr;
ret = hipMemset3D(devPitchedUnPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
SECTION("Reset devPitchedPtr to zero") {
hipPitchedPtr rdevPitchedPtr{};
ret = hipMemset3D(rdevPitchedPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
SECTION("Pass extent fields as max size_t") {
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::max());
ret = hipMemset3D(devPitchedPtr, memsetval, extMW);
REQUIRE(ret != hipSuccess);
ret = hipMemset3D(devPitchedPtr, memsetval, extMH);
REQUIRE(ret != hipSuccess);
if ((TestContext::get()).isAmd()) {
ret = hipMemset3D(devPitchedPtr, memsetval, extMD);
REQUIRE(ret != hipSuccess);
} else {
WARN("Test is skipped for max depth."
<< "Cuda doesn't check the maximum depth of extent field");
}
}
SECTION("Device Ptr out bound and extent(0) passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
hipExtent extent0{};
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0);
// api expected to check extent0 and return success before going for ptr.
REQUIRE(ret == hipSuccess);
}
SECTION("Device Ptr out bound and valid extent passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
HIP_CHECK(hipFree(devPitchedPtr.ptr));
}
TEST_CASE("Unit_hipMemset3DAsync_Negative") {
hipError_t ret;
hipPitchedPtr devPitchedPtr;
hipStream_t stream;
constexpr int memsetval = 1;
constexpr size_t numH = 256;
constexpr size_t numW = 256;
constexpr size_t depth = 10;
constexpr size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
SECTION("Using uninitialized devpitched ptr") {
hipPitchedPtr devPitchedUnPtr;
ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
SECTION("Reset devPitchedPtr to zero") {
hipPitchedPtr rdevPitchedPtr{};
ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
SECTION("Pass extent fields as max size_t") {
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::max());
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream);
REQUIRE(ret != hipSuccess);
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream);
REQUIRE(ret != hipSuccess);
if ((TestContext::get()).isAmd()) {
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream);
REQUIRE(ret != hipSuccess);
} else {
WARN("Test is skipped for max depth."
<< "Cuda doesn't check the maximum depth of extent field");
}
}
SECTION("Device Ptr out bound and extent(0) passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
hipExtent extent0{};
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream);
// api expected to check extent0 and return success before going for ptr.
REQUIRE(ret == hipSuccess);
}
SECTION("Device Ptr out bound and valid extent passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(devPitchedPtr.ptr));
}
projects/hip/tests/catch/unit/memory/hipMemsetInvalidPtr.cc
-127
View File
@@ -1,127 +0,0 @@
/*
* Copyright (c) 2021 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 Scenarios :
1) Test hipMemset apis with invalid pointer and invalid 2D pitch.
2) Test hipMemsetAsync apis with invalid pointer and invalid 2D pitch.
*/
#include <hip_test_common.hh>
#define N 50
#define MEMSETVAL 0x42
/**
* Testcase validates hipMemset apis behavior with
* invalid pointer and invalid 2D pitch value.
*/
TEST_CASE("Unit_hipMemset_InvalidPtrTests") {
hipError_t ret;
constexpr int Nbytes = N*sizeof(char);
char *A_d;
SECTION("hipMemset with null") {
ret = hipMemset(NULL, MEMSETVAL , Nbytes);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemset with hostptr") {
char *A_h;
A_h = reinterpret_cast<char *>(malloc(Nbytes));
ret = hipMemset(A_h, MEMSETVAL, Nbytes);
REQUIRE(ret != hipSuccess);
free(A_h);
}
SECTION("hipMemsetD32 with null") {
ret = hipMemsetD32(NULL, MEMSETVAL , Nbytes);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemsetD16 with null") {
ret = hipMemsetD16(NULL, MEMSETVAL , Nbytes);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemsetD8 with null") {
ret = hipMemsetD8(NULL, MEMSETVAL , Nbytes);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemset2D with null") {
constexpr size_t NUM_H = 256, NUM_W = 256;
size_t pitch_A;
size_t width = NUM_W * sizeof(char);
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
width , NUM_H));
ret = hipMemset2D(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H);
REQUIRE(ret != hipSuccess);
hipFree(A_d);
}
}
/**
* Testcase validates hipMemsetAsync apis behavior with
* invalid pointer and invalid 2D pitch value.
*/
TEST_CASE("Unit_hipMemsetAsync_InvalidPtrTests") {
hipError_t ret;
constexpr int Nbytes = N*sizeof(char);
char *A_d;
SECTION("hipMemsetAsync with null") {
ret = hipMemsetAsync(NULL, MEMSETVAL, Nbytes , 0);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemsetD32Async with null") {
ret = hipMemsetD32Async(NULL, MEMSETVAL , Nbytes, 0);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemsetD16Async with null") {
ret = hipMemsetD16Async(NULL, MEMSETVAL , Nbytes, 0);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemsetD8Async with null") {
ret = hipMemsetD8Async(NULL, MEMSETVAL , Nbytes, 0);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMemset2DAsync with null") {
constexpr size_t NUM_H = 256, NUM_W = 256;
size_t pitch_A;
size_t width = NUM_W * sizeof(char);
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A,
width , NUM_H));
ret = hipMemset2DAsync(NULL, pitch_A, MEMSETVAL, NUM_W, NUM_H, 0);
REQUIRE(ret != hipSuccess);
hipFree(A_d);
}
}
projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc
+229
View File
@@ -0,0 +1,229 @@
/*
* 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 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.
*/
#include <hip_test_common.hh>
#include <limits>
#include "hip/driver_types.h"
static constexpr size_t memsetVal{0x42};
static constexpr hipExtent validExtent{184, 57, 16};
static constexpr size_t height{validExtent.height};
static constexpr size_t width{validExtent.width};
static constexpr hipStream_t nullStream{nullptr};
inline void testHipMemsetApis(void* dst, int value, size_t sizeBytes,
hipError_t expectedReturn = hipErrorInvalidValue) {
HIP_CHECK_ERROR(hipMemset(dst, value, sizeBytes), expectedReturn);
HIP_CHECK_ERROR(hipMemsetAsync(dst, value, sizeBytes, nullStream), expectedReturn);
hipDeviceptr_t devicePtrDst = reinterpret_cast<hipDeviceptr_t>(dst);
HIP_CHECK_ERROR(hipMemsetD32(devicePtrDst, value, sizeBytes), expectedReturn);
HIP_CHECK_ERROR(hipMemsetD32Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
HIP_CHECK_ERROR(hipMemsetD16(devicePtrDst, value, sizeBytes), expectedReturn);
HIP_CHECK_ERROR(hipMemsetD16Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
HIP_CHECK_ERROR(hipMemsetD8(devicePtrDst, value, sizeBytes), expectedReturn);
HIP_CHECK_ERROR(hipMemsetD8Async(devicePtrDst, value, sizeBytes, nullStream), expectedReturn);
}
inline void testHipMemset2DApis(void* dst, size_t pitch, int value, size_t width, size_t height,
hipError_t expectedReturn = hipErrorInvalidValue) {
HIP_CHECK_ERROR(hipMemset2D(dst, pitch, value, width, height), expectedReturn);
HIP_CHECK_ERROR(hipMemset2DAsync(dst, pitch, value, width, height, nullStream), expectedReturn);
}
inline void testHipMemset3DApis(hipPitchedPtr& pitchedDevPtr, int value, const hipExtent& extent,
hipError_t expectedReturn = hipErrorInvalidValue) {
HIP_CHECK_ERROR(hipMemset3D(pitchedDevPtr, value, extent), expectedReturn);
HIP_CHECK_ERROR(hipMemset3DAsync(pitchedDevPtr, value, extent, nullStream), expectedReturn);
}
TEST_CASE("Unit_hipMemset_Negative_InvalidPtr") {
void* dst;
SECTION("Uninitialized Dst") {}
SECTION("Nullptr as dst") { dst = nullptr; }
std::unique_ptr<char[]> hostPtr;
SECTION("Host Pointer as Dst") {
hostPtr.reset(new char[width]);
dst = hostPtr.get();
}
testHipMemsetApis(dst, memsetVal, width);
}
TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsSize") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-20");
#endif
#if !HT_AMD
void* dst;
constexpr size_t outOfBoundsSize{width + 1};
HIP_CHECK(hipMalloc(&dst, width));
testHipMemsetApis(dst, memsetVal, outOfBoundsSize);
HIP_CHECK(hipFree(dst));
#endif
}
TEST_CASE("Unit_hipMemset_Negative_OutOfBoundsPtr") {
void* dst;
HIP_CHECK(hipMalloc(&dst, width));
void* outOfBoundsPtr{reinterpret_cast<char*>(dst) + width + 1};
testHipMemsetApis(outOfBoundsPtr, memsetVal, width);
HIP_CHECK(hipFree(dst));
}
TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") {
void* dst;
SECTION("Uninitialized Dst") {}
SECTION("Nullptr as Dst") { dst = nullptr; }
std::unique_ptr<char[]> hostPtr;
SECTION("Host Pointer as Dst") {
hostPtr.reset(new char[height * width]);
dst = hostPtr.get();
}
void* A_d;
size_t pitch_A;
HIP_CHECK(hipMallocPitch(&A_d, &pitch_A, width, height));
testHipMemset2DApis(dst, pitch_A, memsetVal, width, height);
hipFree(A_d);
}
TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
#endif
void* dst;
size_t realPitch;
HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height));
SECTION("Invalid Pitch") {
size_t invalidPitch = 1;
testHipMemset2DApis(dst, invalidPitch, memsetVal, width, height);
}
SECTION("Invalid Width") {
size_t invalidWidth = realPitch + 1;
testHipMemset2DApis(dst, realPitch, memsetVal, invalidWidth, height);
}
#if !HT_AMD /* EXSWCPHIPT-52 */
SECTION("Invalid height") {
size_t invalidHeight = height + 1;
testHipMemset2DApis(dst, realPitch, memsetVal, width, invalidHeight);
}
#endif
HIP_CHECK(hipFree(dst));
}
TEST_CASE("Unit_hipMemset2D_Negative_OutOfBoundsPtr") {
void* dst;
size_t realPitch;
HIP_CHECK(hipMallocPitch(&dst, &realPitch, width, height));
void* outOfBoundsPtr{reinterpret_cast<char*>(dst) + realPitch * height + 1};
testHipMemset2DApis(outOfBoundsPtr, realPitch, memsetVal, width, height);
HIP_CHECK(hipFree(dst));
}
TEST_CASE("Unit_hipMemset3D_Negative_InvalidPtr") {
hipPitchedPtr pitchedDevPtr;
SECTION("Uninitialized PitchedDevPtr") {}
SECTION("Zero Initialized PitchedDevPtr") { pitchedDevPtr = {}; }
testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent);
}
TEST_CASE("Unit_hipMemset3D_Negative_ModifiedPtr") {
hipPitchedPtr pitchedDevPtr;
HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
void* allocatedMemory{pitchedDevPtr.ptr};
SECTION("Nullptr Dst") { pitchedDevPtr.ptr = nullptr; }
std::unique_ptr<char[]> hostPtr;
SECTION("Host Pointer as Dst") {
hostPtr.reset(new char[validExtent.width * validExtent.height * validExtent.depth]);
pitchedDevPtr.ptr = hostPtr.get();
}
SECTION("Invalid Pitch") { pitchedDevPtr.pitch = 1; }
CAPTURE(pitchedDevPtr.ptr, pitchedDevPtr.pitch, pitchedDevPtr.xsize, pitchedDevPtr.ysize);
testHipMemset3DApis(pitchedDevPtr, memsetVal, validExtent);
HIP_CHECK(hipFree(allocatedMemory));
}
TEST_CASE("Unit_hipMemset3D_Negative_InvalidSizes") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-52");
#endif
hipPitchedPtr pitchedDevPtr;
HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
hipExtent invalidExtent{validExtent};
SECTION("Max Width") { invalidExtent.width = std::numeric_limits<std::size_t>::max(); }
SECTION("Max Height") { invalidExtent.height = std::numeric_limits<std::size_t>::max(); }
#if !HT_NVIDIA /* This case hangs on Nvidia */
SECTION("Max Depth") { invalidExtent.depth = std::numeric_limits<std::size_t>::max(); }
#endif
SECTION("Invalid Width") { invalidExtent.width = pitchedDevPtr.pitch + 1; }
#if !HT_AMD /* EXSWCPHIPT-52 */
SECTION("Invalid height") { invalidExtent.height += 1; }
SECTION("Invalid depth") { invalidExtent.depth += 1; }
#endif
CAPTURE(invalidExtent.width, invalidExtent.height, invalidExtent.depth);
testHipMemset3DApis(pitchedDevPtr, memsetVal, invalidExtent);
HIP_CHECK(hipFree(pitchedDevPtr.ptr));
}
TEST_CASE("Unit_hipMemset3D_Negative_OutOfBounds") {
hipPitchedPtr pitchedDevPtr;
HIP_CHECK(hipMalloc3D(&pitchedDevPtr, validExtent));
hipPitchedPtr outOfBoundsPtr{pitchedDevPtr};
outOfBoundsPtr.ptr = reinterpret_cast<char*>(pitchedDevPtr.ptr) +
pitchedDevPtr.pitch * validExtent.height * validExtent.depth + 1;
SECTION("Extent Equal to 0") {
hipExtent zeroExtent{0, 0, 0};
testHipMemset3DApis(outOfBoundsPtr, memsetVal, zeroExtent, hipSuccess);
}
SECTION("Valid Extent") { testHipMemset3DApis(outOfBoundsPtr, memsetVal, validExtent); }
HIP_CHECK(hipFree(pitchedDevPtr.ptr));
}
projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc
+69
View File
@@ -336,3 +336,72 @@ TEST_CASE("Unit_hipPointerGetAttributes_MultiThread") {
std::thread t4(clusterAllocs, 1000, 1, 4);
if (serialize) t4.join();
}
TEST_CASE("Unit_hipPointerGetAttributes_Negative") {
SECTION("Invalid Attributes Pointer") {
int* dPtr{nullptr};
HIP_CHECK(hipMalloc(&dPtr, sizeof(int)));
HIP_CHECK_ERROR(hipPointerGetAttributes(nullptr, dPtr), hipErrorInvalidValue);
HIP_CHECK(hipFree(dPtr));
}
SECTION("Invalid Device Pointer") {
hipPointerAttribute_t attributes{};
HIP_CHECK_ERROR(hipPointerGetAttributes(&attributes, nullptr), hipErrorInvalidValue);
}
}
// Run this test for all devices for DeviceMemory, HostMemory, ManagedMemory and MappedMemory
TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") {
int deviceCount{0};
HIP_CHECK(hipGetDeviceCount(&deviceCount));
REQUIRE(deviceCount != 0);
// Memory Types
enum MemoryTypes { DeviceMemory = 0, HostMemory, MappedMemory };
auto MemoryType =
GENERATE(MemoryTypes::DeviceMemory, MemoryTypes::HostMemory, MemoryTypes::MappedMemory);
INFO("Memory Type: " << MemoryType);
int* ptr{nullptr};
for (int i = 0; i < deviceCount; i++) {
HIP_CHECK(hipSetDevice(i));
ptr = nullptr;
if (MemoryType == MemoryTypes::DeviceMemory) {
HIP_CHECK(hipMalloc(&ptr, sizeof(int)));
} else if (MemoryType == MemoryTypes::HostMemory) {
HIP_CHECK(hipHostMalloc(&ptr, sizeof(int)));
} else if (MemoryType == MemoryTypes::MappedMemory) {
HIP_CHECK(hipHostMalloc(&ptr, sizeof(int), hipHostMallocMapped));
}
REQUIRE(ptr != nullptr);
hipPointerAttribute_t attributes{};
HIP_CHECK(hipPointerGetAttributes(&attributes, ptr));
REQUIRE(attributes.device == i); // Device Check
// Memory address and type check
if (MemoryType == MemoryTypes::DeviceMemory) {
REQUIRE(attributes.memoryType == hipMemoryTypeDevice);
REQUIRE(attributes.devicePointer == ptr);
REQUIRE(attributes.hostPointer == nullptr);
} else if (MemoryType == MemoryTypes::HostMemory) {
REQUIRE(attributes.memoryType == hipMemoryTypeHost);
REQUIRE(attributes.hostPointer == ptr);
} else if (MemoryType == MemoryTypes::MappedMemory) {
int* devicePtr{nullptr};
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&devicePtr), ptr, 0));
REQUIRE(attributes.hostPointer == ptr);
REQUIRE(attributes.hostPointer == devicePtr);
}
if (MemoryType == MemoryTypes::DeviceMemory) {
HIP_CHECK(hipFree(ptr));
} else if (MemoryType == MemoryTypes::MappedMemory || MemoryType == MemoryTypes::HostMemory) {
HIP_CHECK(hipHostFree(ptr));
}
}
}
projects/hip/tests/catch/unit/stream/hipStreamGetPriority.cc
+52 -34
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -16,59 +16,77 @@ 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.
*/
/**
/*
Testcase Scenarios :
1) Negative tests for hipStreamGetPriority api.
2) Create stream and check default priority of stream is within range.
3) Create stream with high or low priority and check priority is set as expected.
4) Create stream with higher priority or lower priority for the priority range returned.
4) Create stream with higher priority or lower priority for the priority range returned, the stream
priority should be clamped to the priority range.
5) Create stream with CUMask and check priority is returned as expected.
*/
#include <hip_test_common.hh>
/**
* Negative tests for hipStreamGetPriority api.
* Check the error returned when an invalid pointer to a priority is used.
*/
TEST_CASE("Unit_hipStreamGetPriority_Negative") {
hipStream_t stream = 0;
REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
TEST_CASE("Unit_hipStreamGetPriority_InvalidPriorityPointer") {
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK_ERROR(hipStreamGetPriority(stream, nullptr), hipErrorInvalidValue);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with higher priority for the priority range returned.
* Create stream and check priority.
*/
TEST_CASE("Unit_hipStreamGetPriority_higher") {
TEST_CASE("Unit_hipStreamGetPriority_happy") {
int priority_low = 0;
int priority_high = 0;
int devID = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(devID));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
priority_high-1));
hipStream_t stream{};
int priority = 0;
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority == priority_high);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with lower priority for the priority range returned.
*/
TEST_CASE("Unit_hipStreamGetPriority_lower") {
int priority_low = 0;
int priority_high = 0;
int devID = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(devID));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
priority_low+1));
int priority = 0;
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority_low == priority);
HIP_CHECK(hipStreamDestroy(stream));
SECTION("Null Stream") {
HIP_CHECK(hipStreamGetPriority(nullptr, &priority));
// valid priority
REQUIRE(priority_low >= priority);
REQUIRE(priority >= priority_high);
}
SECTION("Created Stream") {
SECTION("Default Priority") {
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamGetPriority(stream, &priority));
// valid priority
// Lower the value higher the priority, higher the value lower the priority
REQUIRE(priority_low >= priority);
REQUIRE(priority >= priority_high);
}
SECTION("High Priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority == priority_high);
}
SECTION("Higher Priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high - 1));
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority == priority_high);
}
SECTION("Low Priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority_low == priority);
}
SECTION("Lower Priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_low + 1));
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority_low == priority);
}
HIP_CHECK(hipStreamDestroy(stream));
}
}
#if HT_AMD
@@ -76,7 +94,7 @@ TEST_CASE("Unit_hipStreamGetPriority_lower") {
* Create stream with CUMask and check priority is returned as expected.
*/
TEST_CASE("Unit_hipStreamGetPriority_StreamsWithCUMask") {
hipStream_t stream;
hipStream_t stream{};
int priority = 0;
int priority_normal = 0;
int priority_low = 0;