Update new tests so as to make them work with new variadic based launch mechanisms.
Этот коммит содержится в:
Alex Voicu
@@ -49,21 +49,39 @@ int main()
|
||||
HIPCHECK(hipMalloc(&Y_d,Nbytes));
|
||||
HIPCHECK(hipMalloc(&Z_d,Nbytes));
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
|
||||
hipLaunchKernel(
|
||||
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));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(1));
|
||||
HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes));
|
||||
HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes));
|
||||
HIPCHECK(hipMemcpyDtoD(Y_d, B_d, Nbytes));
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
|
||||
hipLaunchKernel(
|
||||
HipTest::vectorADD,
|
||||
dim3(blocks),
|
||||
dim3(threadsPerBlock),
|
||||
0,
|
||||
0,
|
||||
static_cast<const int*>(X_d),
|
||||
static_cast<const int*>(Y_d),
|
||||
Z_d,
|
||||
N);
|
||||
HIPCHECK(hipMemcpyDtoH(C_h, Z_d, Nbytes));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
@@ -73,8 +91,8 @@ int main()
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
HIPCHECK(hipFree(Z_d));
|
||||
}
|
||||
|
||||
|
||||
passed();
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
@@ -50,25 +50,43 @@ int main()
|
||||
HIPCHECK(hipMalloc(&Y_d,Nbytes));
|
||||
HIPCHECK(hipMalloc(&Z_d,Nbytes));
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
|
||||
hipLaunchKernel(
|
||||
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));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HIPCHECK(hipStreamCreate(&s));
|
||||
|
||||
HIPCHECK(hipStreamCreate(&s));
|
||||
HIPCHECK(hipSetDevice(1));
|
||||
HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s));
|
||||
HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s));
|
||||
HIPCHECK(hipMemcpyDtoDAsync(Y_d, B_d, Nbytes, s));
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
|
||||
hipLaunchKernel(
|
||||
HipTest::vectorADD,
|
||||
dim3(blocks),
|
||||
dim3(threadsPerBlock),
|
||||
0,
|
||||
0,
|
||||
static_cast<const int*>(X_d),
|
||||
static_cast<const int*>(Y_d),
|
||||
Z_d,
|
||||
N);
|
||||
HIPCHECK(hipMemcpyDtoHAsync(C_h, Z_d, Nbytes, s));
|
||||
HIPCHECK(hipStreamSynchronize(s));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
HIPCHECK(hipStreamDestroy(s));
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
@@ -78,7 +96,7 @@ int main()
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
@@ -48,24 +48,42 @@ int main()
|
||||
HIPCHECK(hipMalloc(&X_d,Nbytes));
|
||||
HIPCHECK(hipMalloc(&Y_d,Nbytes));
|
||||
HIPCHECK(hipMalloc(&Z_d,Nbytes));
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
|
||||
hipLaunchKernel(
|
||||
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));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(1));
|
||||
hipMemcpyPeer(X_d, 1, A_d, 0, Nbytes); //this call is eqv to hipMemcpy(hipMemcpyD2D) which goes via stg bufs.
|
||||
hipMemcpyPeer(Y_d, 1, B_d, 0, Nbytes);
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
|
||||
hipLaunchKernel(
|
||||
HipTest::vectorADD,
|
||||
dim3(blocks),
|
||||
dim3(threadsPerBlock),
|
||||
0,
|
||||
0,
|
||||
static_cast<const int*>(X_d),
|
||||
static_cast<const int*>(Y_d),
|
||||
Z_d,
|
||||
N);
|
||||
HIPCHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipFree(X_d));
|
||||
HIPCHECK(hipFree(Y_d));
|
||||
@@ -74,7 +92,7 @@ int main()
|
||||
passed();
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
@@ -51,26 +51,44 @@ int main()
|
||||
HIPCHECK(hipMalloc(&Y_d,Nbytes));
|
||||
HIPCHECK(hipMalloc(&Z_d,Nbytes));
|
||||
|
||||
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
|
||||
hipLaunchKernel(
|
||||
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));
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HIPCHECK(hipStreamCreate(&s));
|
||||
|
||||
HIPCHECK(hipStreamCreate(&s));
|
||||
HIPCHECK(hipSetDevice(1));
|
||||
HIPCHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, Nbytes, s));
|
||||
HIPCHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, Nbytes, s));
|
||||
|
||||
hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
|
||||
hipLaunchKernel(
|
||||
HipTest::vectorADD,
|
||||
dim3(blocks),
|
||||
dim3(threadsPerBlock),
|
||||
0,
|
||||
0,
|
||||
static_cast<const int*>(X_d),
|
||||
static_cast<const int*>(Y_d),
|
||||
Z_d,
|
||||
N);
|
||||
HIPCHECK ( hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HIPCHECK (hipDeviceSynchronize());
|
||||
HIPCHECK (hipStreamSynchronize(s));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
|
||||
HIPCHECK(hipStreamDestroy(s));
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipFree(X_d));
|
||||
@@ -79,7 +97,7 @@ int main()
|
||||
}
|
||||
|
||||
passed();
|
||||
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
Ссылка в новой задаче
Block a user