In the hipMemset2D and hipMemset3D tests synchronize with the default stream after performing an async memset.

[ROCm/clr commit: cc5abec092]
This commit is contained in:
Vladislav Sytchenko
2019-10-15 17:15:49 -04:00
rodzic 91ceb7dd2b
commit 2bc49fb55c
2 zmienionych plików z 7 dodań i 5 usunięć
@@ -45,7 +45,7 @@ bool testhipMemset2D(int memsetval,int p_gpuDevice)
char *A_d;
char *A_h;
bool testResult = true;
HIPCHECK ( hipMemAllocPitch((hipDeviceptr_t*)&A_d, &pitch_A, width , numH,16) );
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++) {
@@ -89,8 +89,9 @@ bool testhipMemset2DAsync(int memsetval,int p_gpuDevice)
}
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK ( hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
HIPCHECK ( hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream) );
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost));
for (int i=0; i<elements; i++) {
if (A_h[i] != memsetval) {
@@ -53,7 +53,7 @@ bool testhipMemset3D(int memsetval,int p_gpuDevice)
for (size_t i=0; i<elements; i++) {
A_h[i] = 1;
}
HIPCHECK ( hipMemset3D( devPitchedPtr, memsetval, extent) );
HIPCHECK(hipMemset3D( devPitchedPtr, memsetval, extent));
hipMemcpy3DParms myparms = {0};
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
@@ -103,7 +103,8 @@ bool testhipMemset3DAsync(int memsetval,int p_gpuDevice)
}
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK ( hipMemset3DAsync( devPitchedPtr, memsetval, extent, stream) );
HIPCHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
HIPCHECK(hipStreamSynchronize(stream));
hipMemcpy3DParms myparms = {0};
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);