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

Added new Memory API's (#1399)

Browse Source 
Added new memory API's hipMemAllocPitch, hipMemAllocHost, hipMemsetD16, hipMemsetD16Async, hipMemsetD8Async
Modified to support all scenarios hipMemcpyParam2DAsync, hipMemcpyParam2D.
This commit is contained in:
ansurya
2019-10-04 13:36:31 +05:30
committed by Maneesh Gupta
parent d39a2a0749
commit ba9c6e13e4
12 changed files with 404 additions and 48 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/deviceLib/hipTestHalf.cpp
+1 -1
View File
@@ -234,7 +234,7 @@ void checkFunctional() {
int main() {
bool* result{nullptr};
hipHostMalloc(&result, sizeof(result));
hipMemAllocHost((void**)&result, sizeof(result));
result[0] = false;
hipLaunchKernelGGL(
tests/src/runtimeApi/memory/hipArray.cpp
+37 -10
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
@@ -34,6 +34,21 @@ void printSep() {
"======================================================================================\n");
}
inline void initMemCpyParam2D(hip_Memcpy2D &ins, const size_t dpitch,
const size_t spitch, const size_t width,
const size_t height, hipMemoryType dstType,
enum hipMemoryType srcType) {
ins.srcXInBytes=0;
ins.srcY=0;
ins.srcPitch=spitch;
ins.dstXInBytes=0;
ins.dstY=0;
ins.dstPitch=dpitch;
ins.WidthInBytes=width;
ins.Height=height;
ins.dstMemoryType= dstType;
ins.srcMemoryType= srcType;
}
//---
// Test copies of a matrix numW by numH
@@ -65,7 +80,11 @@ void memcpy2Dtest(size_t numW, size_t numH, bool usePinnedHost) {
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW * numH);
HIPCHECK(hipMemcpy2D(A_d, pitch_A, A_h, width, width, numH, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy2D(B_d, pitch_B, B_h, width, width, numH, hipMemcpyHostToDevice));
hip_Memcpy2D ins;
initMemCpyParam2D(ins,pitch_B,width,width,numH,hipMemoryTypeDevice,hipMemoryTypeHost);
ins.dstDevice = (hipDeviceptr_t)B_d;
ins.srcHost = B_h;
HIPCHECK(hipMemcpyParam2D(&ins));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d,
(pitch_C / sizeof(T)) * numH);
@@ -115,7 +134,11 @@ void memcpyArraytest(size_t numW, size_t numH, bool usePinnedHost, bool usePitch
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW * numH);
HIPCHECK(hipMemcpyToArray(A_d, 0, 0, (void*)A_h, width, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpyToArray(B_d, 0, 0, (void*)B_h, width, hipMemcpyHostToDevice));
hip_Memcpy2D ins;
initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeArray,hipMemoryTypeHost);
ins.dstArray = B_d;
ins.srcHost = B_h;
HIPCHECK(hipMemcpyParam2D(&ins));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
(T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW);
@@ -152,15 +175,21 @@ void memcpyArraytest(size_t numW, size_t numH, bool usePinnedHost, bool usePitch
} else {
HIPCHECK(hipMemcpy2DToArray(A_d, 0, 0, (void*)A_h, width, width, numH,
hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy2DToArray(B_d, 0, 0, (void*)B_h, width, width, numH,
hipMemcpyHostToDevice));
hip_Memcpy2D ins;
initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeArray,hipMemoryTypeHost);
ins.dstArray = B_d;
ins.srcHost = B_h;
HIPCHECK(hipMemcpyParam2D(&ins));
}
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
(T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW * numH);
HIPCHECK(hipMemcpy2D((void*)C_h, width, (void*)C_d->data, width, width, numH,
hipMemcpyDeviceToHost));
printf("memcpy srcArray to dstHost\n");
hip_Memcpy2D ins;
initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeHost,hipMemoryTypeArray);
ins.srcArray = C_d;
ins.dstHost = C_h;
HIPCHECK(hipMemcpyParam2D(&ins));
HIPCHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, numW * numH);
@@ -227,8 +256,6 @@ int main(int argc, char* argv[]) {
printf("\n\n=== tests&1 (types)\n");
printSep();
HIPCHECK(hipDeviceReset());
size_t width = N / 6;
size_t height = N / 6;
memcpy2Dtest<float>(321, 211, 0);
memcpy2Dtest<double>(322, 211, 0);
memcpy2Dtest<char>(320, 211, 0);
tests/src/runtimeApi/memory/hipMemset.cpp
+113 -3
View File
@@ -26,11 +26,11 @@ THE SOFTWARE.
* BUILD: %t %s ../../test_common.cpp
* TEST: %t
* //Small copy
* TEST: %t -N 10 --memsetval 0x42 --memsetD32val 0x101
* TEST: %t -N 10 --memsetval 0x42 --memsetD32val 0x101 --memsetD16val 0x10 --memsetD8val 0x1
* // Oddball size
* TEST: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF
* TEST: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF --memsetD16val 0xDEAD --memsetD8val 0xDE
* // Big copy
* TEST: %t -N 256M --memsetval 0xa6 --memsetD32val 0xCAFEBABE
* TEST: %t -N 256M --memsetval 0xa6 --memsetD32val 0xCAFEBABE --memsetD16val 0xCAFE --memsetD8val 0xCA
* HIT_END
*/
@@ -86,6 +86,54 @@ bool testhipMemsetD32(int memsetD32val,int p_gpuDevice)
return testResult;
}
bool testhipMemsetD16(short memsetD16val,int p_gpuDevice)
{
size_t Nbytes = N*sizeof(int);
printf ("testhipMemsetD16 N=%zu memsetD16val=%4x device=%d\n", N, memsetD16val, p_gpuDevice);
short *A_d;
short *A_h;
bool testResult = true;
HIPCHECK ( hipMalloc(&A_d, Nbytes) );
A_h = (short*)malloc(Nbytes);
HIPCHECK ( hipMemsetD16((hipDeviceptr_t)A_d, memsetD16val, N) );
HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
for (int i=0; i<N; i++) {
if (A_h[i] != memsetD16val) {
testResult = false; printf("mismatch at index:%d computed:%08x, memsetD16val:%08x\n", i, A_h[i], memsetD32val);
break;
}
}
HIPCHECK(hipFree(A_d));
free(A_h);
return testResult;
}
bool testhipMemsetD8(char memsetD8val,int p_gpuDevice)
{
size_t Nbytes = N*sizeof(int);
printf ("testhipMemsetD8 N=%zu memsetD8val=%4x device=%d\n", N, memsetD8val, p_gpuDevice);
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMalloc(&A_d, Nbytes) );
A_h = (char*)malloc(Nbytes);
HIPCHECK ( hipMemsetD8((hipDeviceptr_t)A_d, memsetD8val, N) );
HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
for (int i=0; i<N; i++) {
if (A_h[i] != memsetD8val) {
testResult = false; printf("mismatch at index:%d computed:%08x, memsetD8val:%08x\n", i, A_h[i], memsetD8val);
break;
}
}
HIPCHECK(hipFree(A_d));
free(A_h);
return testResult;
}
bool testhipMemsetAsync(int memsetval,int p_gpuDevice)
{
size_t Nbytes = N*sizeof(int);
@@ -144,6 +192,64 @@ bool testhipMemsetD32Async(int memsetD32val,int p_gpuDevice)
return testResult;
}
bool testhipMemsetD16Async(short memsetD16val,int p_gpuDevice)
{
size_t Nbytes = N*sizeof(int);
printf ("testhipMemsetD16Async N=%zu memsetval=%8x device=%d\n", N, memsetD16val, p_gpuDevice);
short *A_d;
short *A_h;
bool testResult = true;
HIPCHECK ( hipMalloc((void**)&A_d, Nbytes) );
A_h = (short*)malloc(Nbytes);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK ( hipMemsetD16Async((hipDeviceptr_t)A_d, memsetD16val, N, stream ));
HIPCHECK ( hipStreamSynchronize(stream));
HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost));
for (int i=0; i<N; i++) {
if (A_h[i] != memsetD16val) {
testResult = false;
printf("mismatch at index:%d computed:%02x, memsetD16val:%02x\n", i, A_h[i], memsetD16val);
break;
}
}
HIPCHECK(hipFree((void*)A_d));
HIPCHECK(hipStreamDestroy(stream));
free(A_h);
return testResult;
}
bool testhipMemsetD8Async(char memsetD8val,int p_gpuDevice)
{
size_t Nbytes = N*sizeof(int);
printf ("testhipMemsetD8Async N=%zu memsetD8val=%2x device=%d\n", N, memsetD8val, p_gpuDevice);
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMalloc((void**)&A_d, Nbytes) );
A_h = (char*)malloc(Nbytes);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK ( hipMemsetD8Async((hipDeviceptr_t)A_d, memsetD8val, N, stream ));
HIPCHECK ( hipStreamSynchronize(stream));
HIPCHECK ( hipMemcpy(A_h, (void*)A_d, Nbytes, hipMemcpyDeviceToHost));
for (int i=0; i<N; i++) {
if (A_h[i] != memsetD8val) {
testResult = false;
printf("mismatch at index:%d computed:%02x, memsetD8val:%02x\n", i, A_h[i], memsetD8val);
break;
}
}
HIPCHECK(hipFree((void*)A_d));
HIPCHECK(hipStreamDestroy(stream));
free(A_h);
return testResult;
}
int main(int argc, char *argv[])
{
HipTest::parseStandardArguments(argc, argv, true);
@@ -153,6 +259,10 @@ int main(int argc, char *argv[])
testResult &= testhipMemsetAsync(memsetval, p_gpuDevice);
testResult &= testhipMemsetD32(memsetD32val, p_gpuDevice);
testResult &= testhipMemsetD32Async(memsetD32val, p_gpuDevice);
testResult &= testhipMemsetD16(memsetD16val, p_gpuDevice);
testResult &= testhipMemsetD16Async(memsetD16val, p_gpuDevice);
testResult &= testhipMemsetD8(memsetD8val, p_gpuDevice);
testResult &= testhipMemsetD8Async(memsetD8val, p_gpuDevice);
if (testResult) passed();
failed("Output Mismatch\n");
}
tests/src/runtimeApi/memory/hipMemset2D.cpp
+8 -5
View File
@@ -45,8 +45,7 @@ bool testhipMemset2D(int memsetval,int p_gpuDevice)
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMallocPitch((void**)&A_d, &pitch_A, width , numH) );
HIPCHECK ( hipMemAllocPitch((hipDeviceptr_t*)&A_d, &pitch_A, width , numH,16) );
A_h = (char*)malloc(sizeElements);
HIPASSERT(A_h != NULL);
for (size_t i=0; i<elements; i++) {
@@ -109,11 +108,15 @@ bool testhipMemset2DAsync(int memsetval,int p_gpuDevice)
int main(int argc, char *argv[])
{
HipTest::parseStandardArguments(argc, argv, true);
bool testResult = false;
HIPCHECK(hipSetDevice(p_gpuDevice));
hipCtx_t context;
hipCtxCreate(&context, 0, p_gpuDevice);
bool testResult = false;
testResult = testhipMemset2D(memsetval, p_gpuDevice);
testResult = testhipMemset2DAsync(memsetval, p_gpuDevice);
passed();
hipCtxDestroy(context);
if(testResult){
passed();
}
}
tests/src/test_common.cpp
+14
View File
@@ -25,6 +25,8 @@ THE SOFTWARE.
size_t N = 4 * 1024 * 1024;
char memsetval = 0x42;
int memsetD32val = 0xDEADBEEF;
short memsetD16val = 0xDEAD;
char memsetD8val = 0xDE;
int iterations = 1;
unsigned blocksPerCU = 6; // to hide latency
unsigned threadsPerBlock = 256;
@@ -106,6 +108,18 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) {
failed("Bad memsetD32val argument");
}
memsetD32val = ex;
} else if (!strcmp(arg, "--memsetD16val")) {
int ex;
if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) {
failed("Bad memsetD16val argument");
}
memsetD16val = ex;
} else if (!strcmp(arg, "--memsetD8val")) {
int ex;
if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) {
failed("Bad memsetD8val argument");
}
memsetD8val = ex;
} else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) {
if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) {
failed("Bad iterations argument");
tests/src/test_common.h
+2
View File
@@ -109,6 +109,8 @@ THE SOFTWARE.
extern size_t N;
extern char memsetval;
extern int memsetD32val;
extern short memsetD16val;
extern char memsetD8val;
extern int iterations;
extern unsigned blocksPerCU;
extern unsigned threadsPerBlock;