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

Use more managed memory interfaces

Browse Source 
Add extra API calls for managed memory: hipMemAdvise, hipMemPrefetchAsync,
hipMemRangeGetAttribute

Change-Id: Ifb09744acdf7582d278453707373c89c30e3ed03
This commit is contained in:
German Andryeyev
2020-10-07 13:40:14 -04:00
parent 543dc8cea5
commit 67c19cb10b
1 changed files with 42 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/runtimeApi/memory/hipMallocManaged.cpp
+42 -8
View File
@@ -3,7 +3,7 @@
#include "test_common.h" #include "test_common.h"
/* HIT_START /* HIT_START
* BUILD: %t %s ../../test_common.cpp * BUILD: %t %s ../../test_common.cpp CLANG_OPTIONS -g -O0
* TEST: %t -N 256M * TEST: %t -N 256M
* HIT_END * HIT_END
*/ */
@@ -20,32 +20,66 @@ void add(int n, float *x, float *y)
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ {
HipTest::parseStandardArguments(argc, argv, true); HipTest::parseStandardArguments(argc, argv, true);
int numElements = N;
printf("info: set device to %d\n", p_gpuDevice);
HIPCHECK(hipSetDevice(p_gpuDevice));
int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N;
bool testResult = true; bool testResult = true;
float *A, *B; float *A, *B;
hipMallocManaged(&A, numElements*sizeof(float)); HIPCHECK(hipMallocManaged(&A, numElements*sizeof(float)));
hipMallocManaged(&B, numElements*sizeof(float)); HIPCHECK(hipMallocManaged(&B, numElements*sizeof(float)));
for (int i = 0; i < numElements; i++) { for (int i = 0; i < numElements; i++) {
A[i] = 1.0f; A[i] = 1.0f;
B[i] = 2.0f; B[i] = 2.0f;
} }
hipDevice_t device = hipCpuDeviceId;
HIPCHECK(hipMemAdvise(A, numElements*sizeof(float), hipMemAdviseSetReadMostly, device));
HIPCHECK(hipMemPrefetchAsync(A, numElements*sizeof(float), 0));
HIPCHECK(hipMemPrefetchAsync(B, numElements*sizeof(float), 0));
HIPCHECK(hipDeviceSynchronize());
HIPCHECK(hipMemRangeGetAttribute(&device, sizeof(device), hipMemRangeAttributeLastPrefetchLocation, A, numElements*sizeof(float)));
if (device != p_gpuDevice) {
printf("hipMemRangeGetAttribute error, device = %d!\n", device);
}
uint32_t read_only = 0xf;
HIPCHECK(hipMemRangeGetAttribute(&read_only, sizeof(read_only), hipMemRangeAttributeReadMostly, A, numElements*sizeof(float)));
if (read_only != 1) {
printf("hipMemRangeGetAttribute error, read_only = %d!\n", read_only);
}
int blockSize = 256; int blockSize = 256;
int numBlocks = (numElements + blockSize - 1) / blockSize; int numBlocks = (numElements + blockSize - 1) / blockSize;
dim3 dimGrid(numBlocks, 1, 1); dim3 dimGrid(numBlocks, 1, 1);
dim3 dimBlock(blockSize, 1, 1); dim3 dimBlock(blockSize, 1, 1);
hipEvent_t event0, event1;
HIPCHECK(hipEventCreate(&event0));
HIPCHECK(hipEventCreate(&event1));
HIPCHECK(hipEventRecord(event0, 0));
hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B); hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B);
HIPCHECK(hipEventRecord(event1, 0));
hipDeviceSynchronize(); HIPCHECK(hipDeviceSynchronize());
float time = 0.0f;
HIPCHECK(hipEventElapsedTime(&time, event0, event1));
printf("Time %.3f ms\n", time);
float maxError = 0.0f; float maxError = 0.0f;
HIPCHECK(hipMemPrefetchAsync(B, numElements*sizeof(float), hipCpuDeviceId));
HIPCHECK(hipDeviceSynchronize());
device = p_gpuDevice;
HIPCHECK(hipMemRangeGetAttribute(&device, sizeof(device), hipMemRangeAttributeLastPrefetchLocation, A, numElements*sizeof(float)));
if (device != hipCpuDeviceId) {
printf("hipMemRangeGetAttribute error (CPU device is expected), device = %d!\n", device);
}
for (int i = 0; i < numElements; i++) for (int i = 0; i < numElements; i++)
maxError = fmax(maxError, fabs(B[i]-3.0f)); maxError = fmax(maxError, fabs(B[i]-3.0f));
hipFree(A); HIPCHECK(hipFree(A));
hipFree(B); HIPCHECK(hipFree(B));
if(maxError == 0.0f) if(maxError == 0.0f)
passed(); passed();
failed("Output Mismatch\n"); failed("Output Mismatch\n");