AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

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

Просмотр исходного кода 
Change-Id: I4202f7a152cc88aeddcd91094a60d5205f809b4f
Этот коммит содержится в:
Jenkins
2022-10-20 07:10:43 -04:00
родитель 80a14899a1 685390469a
Коммит 57ba04d159
22 изменённых файлов: 425 добавлений и 1031 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
include/hip/hip_runtime_api.h
+3 -2
Просмотреть файл
@@ -1707,8 +1707,9 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
* hipErrorInvalidHandle,
* hipErrorTooManyPeers
*
* @note No guarantees are made about the address returned in @p *devPtr.
* In particular, multiple processes may not receive the same address for the same @p handle.
* @note During multiple processes, using the same memory handle opened by the current context,
* there is no guarantee that the same device poiter will be returned in @p *devPtr.
* This is diffrent from CUDA.
*
*/
hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
tests/catch/multiproc/CMakeLists.txt
-1
Просмотреть файл
@@ -9,7 +9,6 @@ set(LINUX_TEST_SRC
hipGetDevicePropertiesMproc.cc
hipSetGetDeviceMproc.cc
hipIpcMemAccessTest.cc
hipHostMallocTestsMproc.cc
hipMallocConcurrencyMproc.cc
hipMemCoherencyTstMProc.cc
hipIpcEventHandle.cc
tests/catch/multiproc/hipHostMallocTestsMproc.cc
-318
Просмотреть файл
@@ -1,318 +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) Validate page table allocation of hipHostMalloc() api when
HIP_VISIBLE_DEVICES set to single device.
2) Validate page table allocation of hipHostMalloc() api when
HIP_VISIBLE_DEVICES set to list of multiple devices.
*/
#include <hip_test_common.hh>
#include <hip_test_smi.hh>
#if defined(__linux__)
#include <unistd.h>
#include <stdlib.h>
#include <sys/wait.h>
#include <dlfcn.h>
#if HT_AMD
/* AMD specific as below tests use rocm-smi */
/**
* Defines
*/
#define LIB_ROCMSMI "librocm_smi64.so"
#define ALLOC_SIZE (30*1024*1024)
/**
* Global variables
*/
static rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t,
rsmi_memory_type_t,
uint64_t *);
static rsmi_status_t (*rsmi_init_fp)(uint64_t);
static rsmi_status_t (*rsmi_shut_down_fp)();
static void *rocm_smi_h;
/**
* Fetches Gpu device count
*/
static void getDeviceCount(int *pdevCnt) {
int fd[2], val = 0;
pid_t childpid;
// create pipe descriptors
pipe(fd);
// disable visible_devices env from shell
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
childpid = fork();
if (childpid > 0) { // Parent
close(fd[1]);
// parent will wait to read the device cnt
read(fd[0], &val, sizeof(val));
// close the read-descriptor
close(fd[0]);
// wait for child exit
wait(nullptr);
*pdevCnt = val;
} else if (!childpid) { // Child
int devCnt = 1;
// writing only, no need for read-descriptor
close(fd[0]);
HIP_CHECK(hipGetDeviceCount(&devCnt));
// send the value on the write-descriptor:
write(fd[1], &devCnt, sizeof(devCnt));
// close the write descriptor:
close(fd[1]);
exit(0);
} else { // failure
*pdevCnt = 0;
return;
}
}
/**
* Initializes rocm smi library handles
*/
static bool rocm_smi_init() {
// Open ROCm SMI Library
if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) {
printf("Error opening rocm smi library!\n");
return false;
}
void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get");
if (!fnsym) {
printf("Error getting rsmi_dev_memory_usage_get() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_dev_memory_usage_get_fp = reinterpret_cast<rsmi_status_t (*)(uint32_t,
rsmi_memory_type_t, uint64_t *)>(fnsym);
fnsym = dlsym(rocm_smi_h, "rsmi_init");
if (!fnsym) {
printf("Error getting rsmi_init() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_init_fp = reinterpret_cast<rsmi_status_t (*)(uint64_t)>(fnsym);
fnsym = dlsym(rocm_smi_h, "rsmi_shut_down");
if (!fnsym) {
printf("Error getting rsmi_shut_down() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_shut_down_fp = reinterpret_cast<rsmi_status_t (*)()>(fnsym);
uint64_t init_flags = 0;
rsmi_status_t retsmi_init;
retsmi_init = rsmi_init_fp(init_flags);
if (RSMI_STATUS_SUCCESS != retsmi_init) {
printf("Error when initializing rocm_smi\n");
dlclose(rocm_smi_h);
return false;
}
return true;
}
/**
* Exits rocm smi library
*/
static void rocm_smi_exit() {
rsmi_shut_down_fp();
dlclose(rocm_smi_h);
}
/**
* Validates page table memory allocations
* by setting visible devices selected.
*/
static bool validatePageTableAllocations(const char *devList, int visDevCnt) {
int fd[2];
bool testResult = false;
pid_t pid;
int numdev = 0;
getDeviceCount(&numdev);
REQUIRE(numdev > 0);
if (pipe(fd) < 0) {
printf("Pipe system call failed\n");
return false;
}
pid = fork();
if (!pid) { // Child process
rsmi_status_t ret;
std::vector<int> prev, current;
uint64_t used = 0;
int tmpdev = 0, changeCnt = 0, indx = 0;
char *ptr = nullptr;
bool testPassed = true;
// Disable visible_devices env from shell
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
setenv("HIP_VISIBLE_DEVICES", devList, 1);
// First Call to initialize hip api
HIP_CHECK(hipGetDeviceCount(&tmpdev));
// Get memory snapshot before hostmalloc
for (indx = 0; indx < numdev; indx++) {
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
if (RSMI_STATUS_SUCCESS != ret) {
printf("Error while running rsmi_dev_memory_usage_get func\n");
dlclose(rocm_smi_h);
rsmi_shut_down_fp();
return false;
}
prev.push_back(used);
}
HIP_CHECK(hipHostMalloc(&ptr, ALLOC_SIZE));
// Get memory snapshot after hostmalloc
for (indx = 0; indx < numdev; indx++) {
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
if (RSMI_STATUS_SUCCESS != ret) {
printf("Error while running rsmi_dev_memory_usage_get func\n");
dlclose(rocm_smi_h);
rsmi_shut_down_fp();
HIP_CHECK(hipHostFree(ptr));
return false;
}
current.push_back(used);
}
for (indx = 0; indx < numdev; indx++) {
// For visible hip devices, there should be increase in VRAM usage
// due to page table allocations
// For NON visible hip devices, there can be reduction in VRAM usage
// due to removal of page tables from them
if (current[indx] > prev[indx])
changeCnt++;
}
// Check if memory allocation happened only for visible devices
if (changeCnt == visDevCnt) {
testPassed = true;
} else {
testPassed = false;
}
HIP_CHECK(hipHostFree(ptr));
// writing only, no need for read-descriptor
close(fd[0]);
// send the value on the write-descriptor:
write(fd[1], &testPassed, sizeof(testPassed));
// close the write descriptor:
close(fd[1]);
exit(0);
} else if (pid > 0) { // parent
close(fd[1]);
read(fd[0], &testResult, sizeof(testResult));
close(fd[0]);
wait(nullptr);
} else {
printf("fork() failed\n");
HIP_ASSERT(false);
}
return testResult;
}
/**
* Test page table allocation when HIP_VISIBLE_DEVICES set to
* single device
*/
TEST_CASE("Unit_hipHostMalloc_SingleVisibleDevicePageAlloc") {
int devCnt;
std::string str;
if (!rocm_smi_init()) {
WARN("Testcase skipped as rocm smi not initialized/present");
return;
}
getDeviceCount(&devCnt);
REQUIRE(devCnt > 0);
// Select single visible device and validate memory usage
for (int i = 0; i < devCnt; i++) {
str = std::to_string(i);
REQUIRE(validatePageTableAllocations(str.c_str(), 1) == true);
}
rocm_smi_exit();
}
/**
* Test page table allocation when HIP_VISIBLE_DEVICES set to
* multiple devices
*/
TEST_CASE("Unit_hipHostMalloc_MultipleVisibleDevicesPageAlloc") {
int devCnt = 0, vdCnt = 0;
std::string str;
if (!rocm_smi_init()) {
WARN("Testcase skipped as rocm smi not initialized/present");
return;
}
getDeviceCount(&devCnt);
REQUIRE(devCnt > 0);
// Select multiple visible devices and validate memory usage
for (int i = 0; i < devCnt; i++) {
if (i == 0)
str += std::to_string(i);
else
str += "," + std::to_string(i);
vdCnt++;
REQUIRE(validatePageTableAllocations(str.c_str(), vdCnt) == true);
}
rocm_smi_exit();
}
#endif // HT_AMD
#endif // __linux__
tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+94 -138
Просмотреть файл
@@ -40,15 +40,12 @@
__global__ void CoherentTst(int *ptr, int PeakClk) {
// Incrementing the value by 1
int64_t GpuFrq = (PeakClk * 1000);
int64_t GpuFrq = int64_t(PeakClk) * 1000;
int64_t StrtTck = clock64();
atomicAdd(ptr, 1);
// The following while loop checks the value in ptr for around 3-4 seconds
while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
if (*ptr == 3) {
atomicAdd(ptr, 1);
return;
}
if (atomicCAS(ptr, 3, 4) == 3) break;
}
}
@@ -57,8 +54,6 @@ __global__ void SquareKrnl(int *ptr) {
*ptr = (*ptr) * (*ptr);
}
// The variable below will work as signal to decide pass/fail
static bool YES_COHERENT = false;
@@ -641,40 +636,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
*Ptr = 4;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
*Ptr = 4;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
@@ -689,40 +675,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
*Ptr = 1;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
*Ptr = 1;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
@@ -736,45 +713,35 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
*Ptr = 4;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
*Ptr = 4;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
/* Test Case Description: The following test checks if the memory exhibits
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
// The following test is AMD specific test hence skipping for Nvidia
@@ -784,42 +751,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
REQUIRE(false);
}
int stat = 0, Pageable = 0;
int stat = 0;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable) {
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
*Ptr = 1;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
if (fork() == 0) { // child process
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
// Allocating hipHostMalloc() memory
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
*Ptr = 1;
TstCoherency(Ptr, HmmMem);
if (YES_COHERENT) {
// exit() with code 10 which indicates pass
HIP_CHECK(hipHostFree(Ptr));
exit(10);
} else {
// exit() with code 9 which indicates fail
HIP_CHECK(hipHostFree(Ptr));
exit(9);
}
} else { // parent process
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
}
}
#endif
tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc
+1 -1
Просмотреть файл
@@ -349,7 +349,7 @@ static bool TestMemoryAllocationInLoop(int test_type,
}
if (!bPassed) break;
}
hipFree(outputVec_d);
HIP_CHECK(hipFree(outputVec_d));
free(outputVec_h);
return bPassed;
}
tests/catch/stress/memory/hipMallocManagedStress.cc
+1 -1
Просмотреть файл
@@ -158,7 +158,7 @@ TEST_CASE("Stress_hipMallocManaged_MultiSize") {
hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm));
dim3 dimBlock(blockSize, 1, 1);
for (int i = 1; i < (1024*1024); ++i) {
for (int i = 1; i < (1024*100); ++i) {
HIP_CHECK(hipMallocManaged(&Hmm1, i));
HIP_CHECK(hipMallocManaged(&Hmm2, i));
for (int j = 0; j < i; ++j) {
tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc
+9 -15
Просмотреть файл
@@ -57,9 +57,8 @@ static int HmmAttrPrint() {
return managed;
}
static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) {
static void ReleaseResource(int *Hmm, hipStream_t *strm) {
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipFree(Hmm1));
HIP_CHECK(hipStreamDestroy(*strm));
}
@@ -70,11 +69,10 @@ static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) {
TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") {
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
int *Hmm = nullptr, *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4);
int *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4);
int InitVal = 123, NumElms = MemSz/4;
bool IfTestPassed = true;
HIP_CHECK(hipGetDeviceCount(&NumDevs));
HIP_CHECK(hipMallocManaged(&Hmm, MemSz));
HIP_CHECK(hipMallocManaged(&Hmm1, MemSz));
for (int i = 0; i < NumElms; ++i) {
Hmm1[i] = InitVal;
@@ -93,44 +91,40 @@ TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") {
// Prefetching memory from i to j
HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, j, strm));
HIP_CHECK(hipStreamSynchronize(strm));
MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms);
MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm1, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
// Verifying the result
for (int m = 0; m < NumElms; ++m) {
if (Hmm[m] != (InitVal * InitVal)) {
if (Hmm1[m] != (InitVal * InitVal)) {
IfTestPassed = false;
}
}
if (!IfTestPassed) {
ReleaseResource(Hmm, Hmm1, &strm);
ReleaseResource(Hmm1, &strm);
INFO("Did not find expected value!");
REQUIRE(false);
}
// Resetting the values in Hmm
HIP_CHECK(hipMemset(Hmm, 0, MemSz));
// Prefetching memory from j to i
HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, strm));
HIP_CHECK(hipStreamSynchronize(strm));
MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms);
MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm1, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
// Verifying the result
for (int m = 0; m < NumElms; ++m) {
if (Hmm[m] != (InitVal * InitVal)) {
if (Hmm1[m] != (InitVal * InitVal)) {
IfTestPassed = false;
}
}
if (!IfTestPassed) {
ReleaseResource(Hmm, Hmm1, &strm);
ReleaseResource(Hmm1, &strm);
INFO("Did not find expected value!");
REQUIRE(false);
}
// Resetting the values in Hmm
HIP_CHECK(hipMemset(Hmm, 0, MemSz));
HIP_CHECK(hipStreamDestroy(strm));
}
}
// Releasing the resources in case all the scenarios passed
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipFree(Hmm1));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
tests/catch/stress/memory/hipMemcpyMThreadMSize.cc
+4 -4
Просмотреть файл
@@ -72,7 +72,7 @@ void Memcpy_And_verify(int NUM_ELM) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIP_CHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDefault));
@@ -122,7 +122,7 @@ void Memcpy_And_verify(int NUM_ELM) {
int canAccessPeer = 0;
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(TestType)));
@@ -165,7 +165,7 @@ void Memcpy_And_verify(int NUM_ELM) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIP_CHECK(hipMemcpyAsync(A_d[j], A_d[i],
NUM_ELM * sizeof(TestType),
@@ -219,7 +219,7 @@ void Memcpy_And_verify(int NUM_ELM) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(j));
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]),
tests/catch/stress/memory/memcpy.cc
+5 -5
Просмотреть файл
@@ -5,25 +5,25 @@ TEST_CASE("Stress_hipMalloc", "DifferentSizes") {
SECTION("Size 10") {
auto res = hipMalloc(&d_a, sizeof(10));
REQUIRE(res == hipSuccess);
hipFree(d_a);
HIP_CHECK(hipFree(d_a));
d_a = nullptr;
}
SECTION("Size 100") {
auto res = hipMalloc(&d_a, sizeof(100));
REQUIRE(res == hipSuccess);
hipFree(d_a);
HIP_CHECK(hipFree(d_a));
d_a = nullptr;
}
SECTION("Size 1000") {
auto res = hipMalloc(&d_a, sizeof(1000));
REQUIRE(res == hipSuccess);
hipFree(d_a);
HIP_CHECK(hipFree(d_a));
d_a = nullptr;
}
SECTION("Size 10000") {
auto res = hipMalloc(&d_a, sizeof(10000));
REQUIRE(res == hipSuccess);
hipFree(d_a);
HIP_CHECK(hipFree(d_a));
d_a = nullptr;
}
SECTION("Size MAX") {
@@ -31,4 +31,4 @@ TEST_CASE("Stress_hipMalloc", "DifferentSizes") {
REQUIRE(res == hipErrorOutOfMemory);
d_a = nullptr;
}
}
}
tests/catch/stress/printf/Stress_printf_ComplexKernels.cc
+1 -1
Просмотреть файл
@@ -494,7 +494,7 @@ TEST_CASE("Stress_printf_ComplexKernelMultStreamMultGpu") {
unsigned int print_limit = 4; // = 4 GB
uint32_t iterCount = 1;
int numOfGPUs = 0;
hipGetDeviceCount(&numOfGPUs);
HIP_CHECK(hipGetDeviceCount(&numOfGPUs));
if (numOfGPUs < 2) {
printf("Skipping test because numOfGPUs < 2\n");
return;
tests/catch/stress/stream/streamEnqueue.cc
+35 -28
Просмотреть файл
@@ -26,11 +26,13 @@ THE SOFTWARE.
#include <random>
#include <thread>
__global__ void addVal(unsigned long long* ptr, size_t index, unsigned long long val) {
__global__ void addVal(unsigned long long* ptr, size_t index,
unsigned long long val) {
atomicAdd(ptr + index, val);
}
// Create a copy constructible AtomicWrap around std::atomic so that we can put it in a vector
// Create a copy constructible AtomicWrap around std::atomic so that
// we can put it in a vector
template <typename T> struct AtomicWrap {
std::atomic<T> data;
@@ -68,18 +70,19 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") {
constexpr size_t maxWork = 10000;
constexpr size_t maxVal = 10;
std::uniform_int_distribution<std::mt19937::result_type> genIndex(0, hwThreads - 1);
std::uniform_int_distribution<std::mt19937::result_type> genIndex(0,
hwThreads - 1);
std::uniform_int_distribution<std::mt19937::result_type> genWork(0, maxWork);
std::uniform_int_distribution<std::mt19937::result_type> genVal(0, maxVal);
auto enqueueKernelThread = [&](hipStream_t stream) {
auto iter = genWork(engine); // Generate work to be done via thread
for (auto i = 0; i < iter; i++) {
for (unsigned long i = 0; i < iter; i++) {
auto index = genIndex(engine); // Generate Index to add to
auto val = genVal(engine); // Generate value to add to the destination
auto val = genVal(engine); // Generate value to add to the destination
hostData[index].data += val; // Replicate it on host
addVal<<<1, 1, 0, stream>>>(dPtr, static_cast<size_t>(index),
static_cast<unsigned long long>(val)); // And on device
static_cast<unsigned long long>(val)); // And on device
}
};
@@ -101,8 +104,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") {
HIP_CHECK(hipStreamDestroy(stream));
auto hPtr = std::make_unique<unsigned long long[]>(hwThreads);
HIP_CHECK(
hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads,
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(dPtr));
@@ -113,7 +116,7 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") {
}
}
__global__ void doOperation(int* dPtr, size_t size, int val) {
__global__ void doOperation(int* dPtr, int val) {
auto i = threadIdx.x;
atomicAdd(dPtr + i, val);
}
@@ -135,14 +138,15 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
std::vector<hipStream_t> streamPool{};
streamPool.reserve(deviceCount * streamPerGPU);
std::map<hipStream_t, int*> streamToDeviceMemory; // Map of stream and device memory
std::map<hipStream_t, AtomicWrap<int>> streamToHostMemory; // Map of stream and host result
std::map<hipStream_t, size_t> streamToDeviceIndex; // Map of stream and device it was created on
// Map of stream and device memory
std::map<hipStream_t, int*> streamToDeviceMemory;
// Map of stream and host result
std::map<hipStream_t, AtomicWrap<int>> streamToHostMemory;
// Map of stream and device it was created on
std::map<hipStream_t, size_t> streamToDeviceIndex;
constexpr size_t size = 1024;
for (size_t i = 0; i < deviceCount; i++) {
for (int i = 0; i < deviceCount; i++) {
HIP_CHECK(hipSetDevice(i));
for (size_t j = 0; j < streamPerGPU; j++) {
@@ -155,8 +159,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
HIP_CHECK(hipMalloc(&dPtr, sizeof(int) * size));
REQUIRE(dPtr != nullptr);
HIP_CHECK(hipMemset(dPtr, 0, sizeof(int) * size));
streamToDeviceMemory[stream] = dPtr; // All streams work on exclusive memory
// All streams work on exclusive memory
streamToDeviceMemory[stream] = dPtr;
streamToHostMemory[stream] = AtomicWrap<int>(0); // CPU result
@@ -171,8 +175,10 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
std::random_device device;
std::mt19937 engine(device());
std::uniform_int_distribution<std::mt19937::result_type> genVal(-maxVal, maxVal);
std::uniform_int_distribution<std::mt19937::result_type> genStream(0, streamPool.size() - 1);
std::uniform_int_distribution<std::mt19937::result_type> genVal(-maxVal,
maxVal);
std::uniform_int_distribution<std::mt19937::result_type> genStream(0,
streamPool.size() - 1);
#if HT_NVIDIA
std::mutex ness; // On nvidia, current device needs to match stream's device
@@ -183,7 +189,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
#if HT_NVIDIA
std::unique_lock<std::mutex> lock(ness); // Lock on creation
#endif
hipStream_t stream = streamPool[genStream(engine)]; // Get a random stream
// Get a random stream
hipStream_t stream = streamPool[genStream(engine)];
// TODO use HIP_CHECK_THREAD when PR#2664 is merged
if (hipSuccess != hipSetDevice(streamToDeviceIndex[stream])) {
@@ -191,11 +198,10 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
}
int val = genVal(engine); // Generate Value to add/sub to
streamToHostMemory[stream].data.fetch_add(val); // Replicate result on CPU
// Replicate result on CPU
streamToHostMemory[stream].data.fetch_add(val);
auto dPtr = streamToDeviceMemory[stream];
doOperation<<<1, 1024, 0, stream>>>(dPtr, size,
val); // On GPU
doOperation<<<1, 1024, 0, stream>>>(dPtr, val); // On GPU
}
};
@@ -219,13 +225,14 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
for (auto& i : streamPool) {
HIP_CHECK(hipStreamSynchronize(i));
auto dResult = std::make_unique<int[]>(size);
HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i], sizeof(int) * size,
hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i],
sizeof(int) * size, hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(streamToDeviceMemory[i]));
HIP_CHECK(hipStreamDestroy(i));
auto res = streamToHostMemory[i].data.load();
INFO("Matching CPU: " << res << " GPU: " << dResult[0] << " Dev Ptr: "
<< streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]);
REQUIRE(std::all_of(dResult.get(), dResult.get() + size, [=](int r) { return r == res; }));
<< streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]);
REQUIRE(std::all_of(dResult.get(), dResult.get() + size,
[=](int r) { return r == res; }));
}
}
tests/catch/unit/CMakeLists.txt
+3
Просмотреть файл
@@ -31,3 +31,6 @@ add_subdirectory(texture)
add_subdirectory(streamperthread)
add_subdirectory(kernel)
add_subdirectory(multiThread)
if(HIP_PLATFORM STREQUAL "amd")
add_subdirectory(clock)
endif()
tests/catch/unit/clock/CMakeLists.txt
+29
Просмотреть файл
@@ -0,0 +1,29 @@
# 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.
# Common Tests - Test independent of all platforms
set(TEST_SRC
hipClockCheck.cc
)
hip_add_exe_to_target(NAME ClockCheckTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
tests/catch/unit/clock/hipClockCheck.cc
+116
Просмотреть файл
@@ -0,0 +1,116 @@
/*
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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip/hip_ext.h>
#define ONESECOND 1000 // in ms
#define HALFSECOND 500 // in ms
enum CLOCK_MODE {
CLOCK_MODE_CLOCK64,
CLOCK_MODE_WALL_CLOCK64
};
__global__ void kernel_c(int clockRate, uint64_t wait_t) {
uint64_t start = clock64() / clockRate, cur = 0; // in ms
do { cur = clock64() / clockRate-start;} while (cur < wait_t);
}
__global__ void kernel_w(int clockRate, uint64_t wait_t) {
uint64_t start = wall_clock64() / clockRate, cur = 0; // in ms
do { cur = wall_clock64() / clockRate-start;} while (cur < wait_t);
}
bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2,
float expectedTime1, float expectedTime2) {
bool testStatus = false;
float ratio = m == CLOCK_MODE_CLOCK64 ? 0.5 : 0.01;
if (fabs(time1 - expectedTime1) < ratio * expectedTime1
&& fabs(time2 - expectedTime2) < ratio * expectedTime2) {
WARN("Succeeded: Expected Vs Actual: Kernel1 - " << expectedTime1 << " Vs " << time1
<< ", Kernel2 - " << expectedTime2 << " Vs " << time2);
testStatus = true;
} else {
FAIL_CHECK("Failed: Expected Vs Actual: Kernel1 -" << expectedTime1 << " Vs " << time1
<< ", Kernel2 - " << expectedTime2 << " Vs " << time2);
testStatus = false;
}
return testStatus;
}
/*
* Launching kernel1 and kernel2 and then we try to
* get the event elapsed time of each kernel using the start and
* end events.The event elapsed time should return us the kernel
* execution time for that particular kernel
*/
bool kernelTimeExecution(CLOCK_MODE m, int clockRate,
uint64_t expectedTime1, uint64_t expectedTime2) {
hipStream_t stream;
hipEvent_t start_event1, end_event1, start_event2, end_event2;
float time1 = 0, time2 = 0;
HIPCHECK(hipEventCreate(&start_event1));
HIPCHECK(hipEventCreate(&end_event1));
HIPCHECK(hipEventCreate(&start_event2));
HIPCHECK(hipEventCreate(&end_event2));
HIPCHECK(hipStreamCreate(&stream));
hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w,
dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, clockRate, expectedTime1);
hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w,
dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, clockRate, expectedTime2);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipEventElapsedTime(&time1, start_event1, end_event1));
HIPCHECK(hipEventElapsedTime(&time2, start_event2, end_event2));
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipEventDestroy(start_event1));
HIPCHECK(hipEventDestroy(end_event1));
HIPCHECK(hipEventDestroy(start_event2));
HIPCHECK(hipEventDestroy(end_event2));
return verifyTimeExecution(m, time1, time2, expectedTime1, expectedTime2);
}
TEST_CASE("Unit_hipClock64_Check") {
HIPCHECK(hipSetDevice(0));
int clockRate = 0; // in KHz
HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeClockRate, 0));
SECTION("Verify kernel execution time via clock64()") {
CHECK(kernelTimeExecution(CLOCK_MODE_CLOCK64, clockRate, ONESECOND, HALFSECOND));
}
}
TEST_CASE("Unit_hipWallClock64_Check") {
HIPCHECK(hipSetDevice(0));
int clockRate = 0; // in KHz
HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeWallClockRate, 0));
if(!clockRate) {
INFO("hipDeviceAttributeWallClockRate has not been supported. Skipped");
return;
}
SECTION("Verify kernel execution time via wall_clock64()") {
CHECK(kernelTimeExecution(CLOCK_MODE_WALL_CLOCK64, clockRate, ONESECOND, HALFSECOND));
}
}
tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc
+4 -19
Просмотреть файл
@@ -43,11 +43,9 @@ Mapping is missing for NVIDIA platform hence skipping the testcases
#include <hip_test_kernels.hh>
constexpr size_t N = 1000000;
#if HT_AMD
/* This test covers the negative scenarios of
hipGraphInstantiateWithFlags API */
TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") {
#if HT_NVIDIA
SECTION("Passing nullptr pGraphExec") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
@@ -67,7 +65,6 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") {
hipGraphExec_t graphExec;
REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess);
}
#endif
}
/*
This function verifies the following scenarios
@@ -160,6 +157,7 @@ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) {
// Instantiate and launch the cloned graph
HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Verify graph execution result
HipTest::checkVectorADD(A_h, B_h, C_h, N);
@@ -248,26 +246,14 @@ by creating dependency graph and instantiate, launching and verifying
the result
*/
TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") {
int numDevices = 0;
int canAccessPeer = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
GraphInstantiateWithFlags_DependencyGraph();
} else {
SUCCEED("Machine does not seem to have P2P");
}
} else {
SUCCEED("skipped the testcase as no of devices is less than 2");
}
GraphInstantiateWithFlags_DependencyGraph();
}
/*
This testcase verifies hipGraphInstantiateWithFlags API
by creating dependency graph on GPU-0 and instantiate, launching and verifying
the result on GPU-1
*/
#if HT_NVIDIA
TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") {
int numDevices = 0;
int canAccessPeer = 0;
@@ -283,7 +269,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") {
SUCCEED("skipped the testcase as no of devices is less than 2");
}
}
#endif
/*
This testcase verifies hipGraphInstantiateWithFlags API
by creating capture graph and instantiate, launching and verifying
@@ -325,4 +311,3 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") {
SUCCEED("skipped the testcase as no of devices is less than 2");
}
}
#endif
tests/catch/unit/memory/hipMemAdvise.cc
+54 -89
Просмотреть файл
@@ -664,98 +664,63 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
REQUIRE(false);
}
// The following code block is used to check for gfx906/8 so as to skip if
// any of the gpus available
int fd1[2]; // Used to store two ends of first pipe
pid_t p;
if (pipe(fd1) == -1) {
fprintf(stderr, "Pipe Failed");
REQUIRE(false);
}
// The following code block checks for gfx90a so as to skip if the device is not MI200
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
int GpuId[2] = {0, 0};
p = fork();
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if (p < 0) {
fprintf(stderr, "fork Failed");
REQUIRE(false);
} else if (p > 0) { // parent process
close(fd1[1]); // Close writing end of first pipe
// Wait for child to send a string
wait(NULL);
// Read string from child and close reading end.
read(fd1[0], GpuId, 2 * sizeof(int));
close(fd1[0]);
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
WARN("This test is not applicable on MI60 & MI100."
"Skipping the test!!");
exit(0);
}
} else { // child process
close(fd1[0]); // Close read end of first pipe
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
char *p = NULL;
p = strstr(prop.gcnArchName, "gfx906");
if (p) {
WARN("gfx906 gpu found on this system!!");
GpuId[0] = 1;
}
p = strstr(prop.gcnArchName, "gfx908");
if (p) {
WARN("gfx908 gpu found on this system!!");
GpuId[1] = 1;
}
// Write concatenated string and close writing end
write(fd1[1], GpuId, 2 * sizeof(int));
close(fd1[1]);
exit(0);
}
int stat = 0;
if (fork() == 0) {
// The below part should be inside fork
int managed = HmmAttrPrint();
if (managed == 1) {
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096*4));
for (int i = 0; i < NumElms; ++i) {
Mllc[i] = InitVal;
}
hipStream_t strm;
int DataMismatch = 0;
HIP_CHECK(hipStreamCreate(&strm));
// The following hipMemAdvise() call is made to know if advise on
// aligned_alloc() is causing any issue
HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
HIP_CHECK(hipStreamSynchronize(strm));
MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
for (int i = 0; i < NumElms; ++i) {
if (Mllc[i] != (InitVal + 10)) {
DataMismatch++;
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
int stat = 0;
if (fork() == 0) {
// The below part should be inside fork
int managed = HmmAttrPrint();
if (managed == 1) {
int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123;
// Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz));
Mllc = reinterpret_cast<int*>(aligned_alloc(4096, 4096*4));
for (int i = 0; i < NumElms; ++i) {
Mllc[i] = InitVal;
}
hipStream_t strm;
int DataMismatch = 0;
HIP_CHECK(hipStreamCreate(&strm));
// The following hipMemAdvise() call is made to know if advise on
// aligned_alloc() is causing any issue
HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
HIP_CHECK(hipStreamSynchronize(strm));
MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
HIP_CHECK(hipStreamSynchronize(strm));
for (int i = 0; i < NumElms; ++i) {
if (Mllc[i] != (InitVal + 10)) {
DataMismatch++;
}
}
if (DataMismatch != 0) {
WARN("DataMismatch observed!!");
exit(9); // 9 for failure
} else {
exit(10); // 10 for Pass result
}
}
if (DataMismatch != 0) {
WARN("DataMismatch observed!!");
exit(9); // 9 for failure
} else {
exit(10); // 10 for Pass result
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
} else {
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
wait(&stat);
int Result = WEXITSTATUS(stat);
if (Result != 10) {
REQUIRE(false);
}
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
#endif
@@ -777,7 +742,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -809,7 +774,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
} else {
@@ -847,7 +812,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") {
}
HIP_CHECK(hipFree(Hmm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -929,7 +894,7 @@ TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") {
HIP_CHECK(hipFree(Hmm));
HIP_CHECK(hipStreamDestroy(strm));
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
@@ -955,7 +920,7 @@ TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") {
REQUIRE(false);
}
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
"attribute. Hence skipping the testing with Pass result.\n");
}
}
tests/catch/unit/memory/hipMemAdviseMmap.cc
+1 -1
Просмотреть файл
@@ -50,7 +50,7 @@ TEST_CASE("Unit_hipMemAdvise_MmapMem") {
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if ((managed == 1) && (PageableMem == 1)) {
if (PageableMem == 1) {
#ifdef __linux__
// For now this test is enabled only for linux platforms
FILE *fptr;
tests/catch/unit/memory/hipMemCoherencyTst.cc
+17 -31
Просмотреть файл
@@ -89,33 +89,24 @@ static void TstCoherency(int* ptr, bool hmmMem) {
// passing
#if HT_AMD
TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0;
int *Ptr = nullptr, SIZE = sizeof(int);
bool HmmMem = false;
YES_COHERENT = false;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
if (Pageable == 1) {
// Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
}
SECTION("hipHostMalloc with Default flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
}
SECTION("hipHostMalloc with hipHostMallocMapped flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
}
TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr));
REQUIRE(YES_COHERENT);
} else {
SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
"attribute. Hence skipping the test with Pass result.\n");
// Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
}
SECTION("hipHostMalloc with Default flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
}
SECTION("hipHostMalloc with hipHostMallocMapped flag") {
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
}
TstCoherency(Ptr, HmmMem);
HIP_CHECK(hipHostFree(Ptr));
REQUIRE(YES_COHERENT);
}
#endif
@@ -126,19 +117,14 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
// passing
#if HT_AMD
TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0;
int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
bool HmmMem = true;
YES_COHERENT = false;
HIP_CHECK(hipDeviceGetAttribute(&Pageable,
hipDeviceAttributePageableMemoryAccess, 0));
INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if (managed == 1 && Pageable == 1) {
if (managed == 1) {
// Allocating hipMallocManaged() memory
SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
@@ -150,7 +136,7 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
HIP_CHECK(hipFree(Ptr));
REQUIRE(YES_COHERENT);
} else {
SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
SUCCEED("GPU 0 doesn't support ManagedMemory "
"device attribute. Hence skipping the test with Pass result.\n");
}
}
tests/catch/unit/memory/hipMemcpyPeer.cc
+4 -2
Просмотреть файл
@@ -73,11 +73,13 @@ TEST_CASE("Unit_hipMemcpyPeer_Negative") {
}
SECTION("Passing invalid Destination device ID") {
REQUIRE(hipMemcpyPeer(B_d, 10, A_d, 0, copy_bytes) != hipSuccess);
REQUIRE(hipMemcpyPeer(B_d, numDevices, A_d, 0, copy_bytes) !=
hipSuccess);
}
SECTION("Passing invalid Source device ID") {
REQUIRE(hipMemcpyPeer(B_d, 1, A_d, 10, copy_bytes) != hipSuccess);
REQUIRE(hipMemcpyPeer(B_d, 1, A_d, numDevices, copy_bytes) !=
hipSuccess);
}
HipTest::freeArrays<int>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
} else {
tests/catch/unit/memory/hipMemcpyPeerAsync.cc
+4 -4
Просмотреть файл
@@ -82,12 +82,12 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Negative") {
}
SECTION("Passing invalid Destination device ID") {
REQUIRE(hipMemcpyPeerAsync(B_d, 10, A_d, 0, copy_bytes,
REQUIRE(hipMemcpyPeerAsync(B_d, numDevices, A_d, 0, copy_bytes,
stream) != hipSuccess);
}
SECTION("Passing invalid Source device ID") {
REQUIRE(hipMemcpyPeerAsync(B_d, 10, A_d, 0, copy_bytes,
REQUIRE(hipMemcpyPeerAsync(B_d, 0, A_d, numDevices, copy_bytes,
stream) != hipSuccess);
}
@@ -155,14 +155,14 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
stream));
HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") {
HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes,
hipStreamPerThread));
HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes,
hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(X_d),
tests/src/runtimeApi/memory/hipHostMallocTests.cpp
+35 -366
Просмотреть файл
@@ -1,5 +1,5 @@
/*
Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2020 - 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
@@ -28,323 +28,16 @@ Testcase Scenarios :
4) Pass size as zero for hipHostMalloc() api and check ptr is reset with
with return value success.
(TestCase 2)::
5) Validate memory usage of hipHostMalloc() api when HIP_VISIBLE_DEVICES set
to single device.
6) Validate memory usage of hipHostMalloc() api when HIP_VISIBLE_DEVICES set
to list of multiple devices.
*/
/* Tests 2 and 3 are rocclr specific tests and not supported on nvidia */
/* HIT_START
* BUILD_CMD: %t %hc %S/%s %S/../../test_common.cpp -I%S/../../ -o %T/%t -ldl -std=c++11
* TEST: %t --tests 1
* TEST: %t --tests 2
* TEST: %t --tests 3
* BUILD_CMD: %t %hc %S/%s %S/../../test_common.cpp -I%S/../../ -o %T/%t -std=c++11
* TEST: %t
* HIT_END
*/
#ifdef __linux__
#include <unistd.h>
#include <stdlib.h>
#include <sys/wait.h>
#include <dlfcn.h>
#endif
#include <string>
#include <vector>
#include <limits>
#include "hipHostMallocTests.h"
/**
* Defines
*/
#define LIB_ROCMSMI "librocm_smi64.so"
#include "test_common.h"
#define NUM_BYTES 1000
#define ALLOC_SIZE (30*1024*1024)
/**
* Global variables
*/
rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t, rsmi_memory_type_t,
uint64_t *);
rsmi_status_t (*rsmi_init_fp)(uint64_t);
rsmi_status_t (*rsmi_shut_down_fp)();
void *rocm_smi_h;
/**
* Fetches Gpu device count
*/
void getDeviceCount(int *pdevCnt) {
#ifdef __linux__
int fd[2], val = 0;
pid_t childpid;
// create pipe descriptors
pipe(fd);
// disable visible_devices env from shell
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
childpid = fork();
if (childpid > 0) { // Parent
close(fd[1]);
// parent will wait to read the device cnt
read(fd[0], &val, sizeof(val));
// close the read-descriptor
close(fd[0]);
// wait for child exit
wait(NULL);
*pdevCnt = val;
} else if (!childpid) { // Child
int devCnt = 1;
// writing only, no need for read-descriptor
close(fd[0]);
HIPCHECK(hipGetDeviceCount(&devCnt));
// send the value on the write-descriptor:
write(fd[1], &devCnt, sizeof(devCnt));
// close the write descriptor:
close(fd[1]);
exit(0);
} else { // failure
*pdevCnt = 1;
return;
}
#else
HIPCHECK(hipGetDeviceCount(pdevCnt));
#endif
}
#if defined(__linux__)
/**
* Initializes rocm smi library handles
*/
bool rocm_smi_init() {
// Open ROCm SMI Library
if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) {
printf("Error opening rocm smi library!\n");
return false;
}
void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get");
if (!fnsym) {
printf("Error getting rsmi_dev_memory_usage_get() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_dev_memory_usage_get_fp = reinterpret_cast<rsmi_status_t (*)(uint32_t,
rsmi_memory_type_t, uint64_t *)>(fnsym);
fnsym = dlsym(rocm_smi_h, "rsmi_init");
if (!fnsym) {
printf("Error getting rsmi_init() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_init_fp = reinterpret_cast<rsmi_status_t (*)(uint64_t)>(fnsym);
fnsym = dlsym(rocm_smi_h, "rsmi_shut_down");
if (!fnsym) {
printf("Error getting rsmi_shut_down() function\n");
dlclose(rocm_smi_h);
return false;
}
rsmi_shut_down_fp = reinterpret_cast<rsmi_status_t (*)()>(fnsym);
uint64_t init_flags = 0;
rsmi_status_t retsmi_init;
retsmi_init = rsmi_init_fp(init_flags);
if (RSMI_STATUS_SUCCESS != retsmi_init) {
printf("Error when initializing rocm_smi\n");
dlclose(rocm_smi_h);
return false;
}
return true;
}
/**
* Exits rocm smi library
*/
void rocm_smi_exit() {
rsmi_shut_down_fp();
dlclose(rocm_smi_h);
}
/**
* Validates page table memory allocations
* by setting visible devices selected.
*/
bool validatePageTableAllocations(const char *devList, int devCnt) {
int fd[2];
bool testResult = false;
pid_t pid;
int numdev = 0;
getDeviceCount(&numdev);
if (pipe(fd) < 0) {
printf("Pipe system call failed\n");
return false;
}
pid = fork();
if (!pid) { // Child process
rsmi_status_t ret;
std::vector<int> prev, current;
uint64_t used = 0;
int tmpdev = 0, changeCnt = 0, indx = 0;
char *ptr = NULL;
bool testPassed = true;
// Disable visible_devices env from shell
unsetenv("ROCR_VISIBLE_DEVICES");
unsetenv("HIP_VISIBLE_DEVICES");
setenv("HIP_VISIBLE_DEVICES", devList, 1);
// First Call to initialize hip api
hipGetDeviceCount(&tmpdev);
// Get memory snapshot before hostmalloc
for (indx = 0; indx < numdev; indx++) {
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
if (RSMI_STATUS_SUCCESS != ret) {
printf("Error while running rsmi_dev_memory_usage_get func\n");
dlclose(rocm_smi_h);
rsmi_shut_down_fp();
return false;
}
prev.push_back(used);
}
HIPCHECK(hipHostMalloc(&ptr, ALLOC_SIZE));
// Get memory snapshot after hostmalloc
for (indx = 0; indx < numdev; indx++) {
ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used);
if (RSMI_STATUS_SUCCESS != ret) {
printf("Error while running rsmi_dev_memory_usage_get func\n");
dlclose(rocm_smi_h);
rsmi_shut_down_fp();
hipHostFree(ptr);
return false;
}
current.push_back(used);
}
for (indx = 0; indx < numdev; indx++) {
// For visible hip devices, there should be increase in VRAM usage
// due to page table allocations
// For NON visible hip devices, there can be reduction in VRAM usage
// due to removal of page tables from them
if (current[indx] > prev[indx])
changeCnt++;
}
// Check if memory allocation happened only for visible devices
if (changeCnt == devCnt) {
testPassed = true;
} else {
testPassed = false;
}
hipHostFree(ptr);
// writing only, no need for read-descriptor
close(fd[0]);
// send the value on the write-descriptor:
write(fd[1], &testPassed, sizeof(testPassed));
// close the write descriptor:
close(fd[1]);
exit(0);
} else if (pid > 0) { // parent
close(fd[1]);
read(fd[0], &testResult, sizeof(testResult));
close(fd[0]);
wait(NULL);
} else {
printf("fork() failed\n");
testResult = false;
}
return testResult;
}
/**
* Validate memory usage selecting single visible device
*/
bool validateHostMallocSingleVisibleDevice() {
int devCnt;
std::string str;
bool TestPassed = true;
if (!rocm_smi_init()) {
printf("%s Testcase skipped as rocm smi not initialized/present\n",
__func__);
return true;
}
getDeviceCount(&devCnt);
// Select single visible device and validate memory usage
for (int i = 0; i < devCnt; i++) {
str = std::to_string(i);
TestPassed = validatePageTableAllocations(str.c_str(), 1);
if (!TestPassed)
break;
}
rocm_smi_exit();
return TestPassed;
}
/**
* Validate memory usage selecting multiple visible devices
*/
bool validateHostMallocMultipleVisibleDevices() {
int devCnt = 0, vdCnt = 0;
std::string str;
bool TestPassed = true;
if (!rocm_smi_init()) {
printf("%s Testcase skipped as rocm smi not initialized/present\n",
__func__);
return true;
}
getDeviceCount(&devCnt);
// Select multiple visible devices and validate memory usage
for (int i = 0; i < devCnt; i++) {
if (i == 0)
str += std::to_string(i);
else
str += "," + std::to_string(i);
vdCnt++;
TestPassed = validatePageTableAllocations(str.c_str(), vdCnt);
if (!TestPassed)
break;
}
rocm_smi_exit();
return TestPassed;
}
#endif
int main(int argc, char *argv[]) {
HipTest::parseStandardArguments(argc, argv, true);
@@ -353,63 +46,40 @@ int main(int argc, char *argv[]) {
size_t allocSize = NUM_BYTES;
char *ptr;
if (p_tests == 1) {
// Pass ptr as nullptr.
if ((ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize))
!= hipErrorInvalidValue) {
printf("ArgValidation : Inappropritate error value returned for "
"ptr as nullptr. Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Pass ptr as nullptr.
if ((ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize))
!= hipErrorInvalidValue) {
printf("ArgValidation : Inappropritate error value returned for "
"ptr as nullptr. Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Size as max(size_t).
if ((ret = hipHostMalloc(&ptr,
std::numeric_limits<std::size_t>::max()))
!= hipErrorOutOfMemory) {
printf("ArgValidation : Inappropritate error value returned for "
"max(size_t). Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Size as max(size_t).
if ((ret = hipHostMalloc(&ptr,
std::numeric_limits<std::size_t>::max()))
!= hipErrorOutOfMemory) {
printf("ArgValidation : Inappropritate error value returned for "
"max(size_t). Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Flags as max(uint).
if ((ret = hipHostMalloc(&ptr, allocSize,
std::numeric_limits<unsigned int>::max()))
!= hipErrorInvalidValue) {
printf("ArgValidation : Inappropritate error value returned for "
"max(uint). Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Flags as max(uint).
if ((ret = hipHostMalloc(&ptr, allocSize,
std::numeric_limits<unsigned int>::max()))
!= hipErrorInvalidValue) {
printf("ArgValidation : Inappropritate error value returned for "
"max(uint). Error: '%s'(%d)\n",
hipGetErrorString(ret), ret);
TestPassed &= false;
}
// Pass size as zero and check ptr reset.
HIPCHECK(hipHostMalloc(&ptr, 0));
if (ptr) {
TestPassed &= false;
printf("ArgValidation : ptr is not reset when size(0)\n");
}
} else if (p_tests == 2) {
// Test page table allocation when HIP_VISIBLE_DEVICES set to
// single device
#if defined(__linux__)
TestPassed = validateHostMallocSingleVisibleDevice();
#else
printf("Test validateHostMallocSingleVisibleDevice skipped on"
"non-linux\n");
#endif
} else if (p_tests == 3) {
// Test page table allocation when HIP_VISIBLE_DEVICES set to
// multiple devices
#if defined(__linux__)
TestPassed = validateHostMallocMultipleVisibleDevices();
#else
printf("Test validateHostMallocMultipleVisibleDevices skipped on"
"non-linux\n");
#endif
} else {
printf("Didnt receive any valid option. Try options 1 to 3\n");
TestPassed = false;
// Pass size as zero and check ptr reset.
HIPCHECK(hipHostMalloc(&ptr, 0));
if (ptr) {
TestPassed &= false;
printf("ArgValidation : ptr is not reset when size(0)\n");
}
if (TestPassed) {
@@ -418,4 +88,3 @@ int main(int argc, char *argv[]) {
failed("hipHostMallocTests validation Failed!");
}
}
tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp
+5 -5
Просмотреть файл
@@ -34,8 +34,8 @@
#include "test_common.h"
#include "hip/hip_ext.h"
#define FOURSEC_KERNEL 4999
#define TWOSEC_KERNEL 2999
#define FIVESEC_KERNEL 4999
#define THREESEC_KERNEL 2999
__device__ int globalvar = 1;
__global__ void TwoSecKernel(int clockrate) {
@@ -163,12 +163,12 @@ bool KernelTimeExecution() {
e = hipEventElapsedTime(&time_4sec, start_event1, end_event1);
e = hipEventElapsedTime(&time_2sec, start_event2, end_event2);
if ( (time_4sec < static_cast<float>(FOURSEC_KERNEL)) &&
(time_2sec < static_cast<float>(TWOSEC_KERNEL))) {
if ( (time_4sec < static_cast<float>(FIVESEC_KERNEL)) &&
(time_2sec < static_cast<float>(THREESEC_KERNEL))) {
testStatus = true;
} else {
printf("Expected Vs Actual: Kernel1-<%d Vs %f Kernel2-<%d Vs %f\n",
FOURSEC_KERNEL, time_4sec, TWOSEC_KERNEL, time_2sec);
FIVESEC_KERNEL, time_4sec, THREESEC_KERNEL, time_2sec);
testStatus = false;
}