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

Catch2 Test Fixes (#2961)

Browse Source 
* test fixes

* address PR comment

* PR comment fixing Nvidia pass

* add additional kernel launch checks

* pr comments


[ROCm/hip-tests commit: da453211ed]
This commit is contained in:
Paulius Velesko
2022-10-01 01:07:25 +03:00
committed by GitHub
parent c9c04d65da
commit 65a428aaff
94 changed files with 177 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/memory/hipHostRegister.cc
+2 -1
View File
@@ -101,7 +101,8 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", i
// Reference the registered device pointer Ad from inside the kernel:
for (int i = 0; i < num_devices; i++) {
HIP_CHECK(hipSetDevice(i));
hipLaunchKernelGGL(Inc, dim3(LEN / 512), dim3(512), 0, 0, Ad[i]);
hipLaunchKernelGGL(Inc, dim3(LEN / 32), dim3(32), 0, 0, Ad[i]);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
}
REQUIRE(A[10] == 1 + static_cast<TestType>(num_devices));
projects/hip-tests/catch/unit/memory/hipMallocConcurrency.cc
+2 -2
View File
@@ -111,7 +111,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
@@ -219,7 +219,7 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
projects/hip-tests/catch/unit/memory/hipMallocManaged.cc
+1
View File
@@ -114,6 +114,7 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") {
HIP_CHECK(hipEventRecord(event0, 0));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const float*>(A), static_cast<const float*>(B), C, numElements);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipEventRecord(event1, 0));
HIP_CHECK(hipDeviceSynchronize());
float time = 0.0f;
projects/hip-tests/catch/unit/memory/hipMallocPitch.cc
+1
View File
@@ -536,6 +536,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", ""
hipLaunchKernelGGL(copy_var<TestType>, dim3(1), dim3(1),
0, 0, static_cast<TestType*>(A_d),
static_cast<TestType*>(B_d), ROWS, pitch_A);
HIP_CHECK(hipGetLastError());
// hipMemcpy2D Device to Host
projects/hip-tests/catch/unit/memory/hipMemcpy.cc
+4 -3
View File
@@ -241,6 +241,7 @@ void memcpytest2(DeviceMemory<T>* dmem, HostMemory<T>* hmem,
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const T*>(dmem->A_d()), static_cast<const T*>(dmem->B_d()),
dmem->C_d(), numElements);
HIP_CHECK(hipGetLastError());
if (useDeviceToDevice) {
// Do an extra device-to-device copy here to mix things up:
@@ -402,7 +403,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_KernelLaunch", "", int, float,
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
@@ -567,7 +568,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch",
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
@@ -595,7 +596,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch",
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
projects/hip-tests/catch/unit/memory/hipMemcpyAllApiNegative.cc
+1
View File
@@ -319,6 +319,7 @@ TEST_CASE("Unit_hipMemcpy_HalfMemCopy") {
HIP_CHECK(hipMemcpyAsync(B_h, A_d,
(NUM_ELM/2)*sizeof(float),
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
HipTest::freeArrays<float>(A_d, B_d, C_d, A_h, B_h, C_h, false);
projects/hip-tests/catch/unit/memory/hipMemcpyAsync.cc
+5 -3
View File
@@ -50,6 +50,7 @@ void Thread_func(T *A_d, T *B_d, T* C_d, T* C_h, size_t Nbytes,
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0,
mystream, A_d, C_d, N_ELMTS);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following two MemcpyAsync calls are for sole
// purpose of loading stream with multiple async calls
@@ -76,6 +77,7 @@ void Thread_func_MultiStream() {
hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks),
dim3(threadsPerBlock), 0,
mystream, A_d, C_d, N_ELMTS);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following hipMemcpyAsync() is called only to
// load stream with multiple Async calls
@@ -121,7 +123,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float,
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
@@ -348,7 +350,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch",
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
@@ -379,7 +381,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch",
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyAsync(C_h, Z_d, Nbytes,
hipMemcpyDeviceToHost, gpu1Stream));
HIP_CHECK(hipStreamSynchronize(gpu1Stream));
projects/hip-tests/catch/unit/memory/hipMemcpyDtoD.cc
+2
View File
@@ -66,6 +66,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "",
dim3(1), 0, 0,
static_cast<const TestType *>(A_d),
static_cast<const TestType *>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
@@ -80,6 +81,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "",
dim3(1), 0, 0,
static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
projects/hip-tests/catch/unit/memory/hipMemcpyDtoDAsync.cc
+2
View File
@@ -68,6 +68,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic", "",
dim3(1), 0, 0,
static_cast<const TestType *>(A_d),
static_cast<const TestType *>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
@@ -84,6 +85,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic", "",
dim3(1), 0, 0,
static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpyDtoHAsync(C_h, (hipDeviceptr_t)Z_d, Nbytes, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipDeviceSynchronize());
projects/hip-tests/catch/unit/memory/hipMemcpyPeer.cc
+2
View File
@@ -127,6 +127,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Basic") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
@@ -138,6 +139,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Basic") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(X_d),
static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
projects/hip-tests/catch/unit/memory/hipMemcpyPeerAsync.cc
+4
View File
@@ -143,6 +143,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
@@ -166,6 +167,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(X_d),
static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
@@ -227,6 +229,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_StreamOnDiffDevice") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, C_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
HipTest::checkVectorADD<int>(A_h, B_h, C_h, numElements);
@@ -240,6 +243,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_StreamOnDiffDevice") {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1),
0, 0, static_cast<const int*>(X_d),
static_cast<const int*>(Y_d), Z_d, numElements*sizeof(int));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(C_h, Z_d, numElements*sizeof(int),
hipMemcpyDeviceToHost));
projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc
+8
View File
@@ -99,6 +99,7 @@ void TestwithOnestream(void) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
@@ -136,6 +137,7 @@ void TestwithTwoStream(void) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NUM_STREAMS; ++i) {
@@ -191,6 +193,7 @@ void TestDtoDonSameDevice(void) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NUM_STREAMS; ++i) {
@@ -262,6 +265,7 @@ void TestOnMultiGPUwithOneStream(void) {
dim3(threadsPerBlock), 0, stream[i],
static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NumDevices; ++i) {
@@ -298,6 +302,7 @@ void TestkindDtoH(void) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
hipMemcpyDeviceToHost, stream));
@@ -365,6 +370,7 @@ void TestkindDtoD(void) {
dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NumDevices; ++i) {
@@ -411,6 +417,7 @@ void TestkindDefault(void) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
@@ -472,6 +479,7 @@ void TestkindDefaultForDtoD(void) {
dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NumDevices; ++i) {
projects/hip-tests/catch/unit/memory/hipMemcpyWithStreamMultiThread.cc
+8
View File
@@ -167,6 +167,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool &val_res) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h, B_h, C_h);
@@ -203,6 +204,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool &val_res) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NoofStreams; ++i) {
@@ -258,6 +260,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < NoofStreams; ++i) {
@@ -329,6 +332,7 @@ void HipMemcpyWithStreamMultiThreadtests::
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < numDevices; ++i) {
@@ -363,6 +367,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool &val_res) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
hipMemcpyDeviceToHost, stream));
@@ -434,6 +439,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < numDevices; ++i) {
@@ -481,6 +487,7 @@ void HipMemcpyWithStreamMultiThreadtests::
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
val_res = ValidateResult(A_h, B_h, C_h);
@@ -544,6 +551,7 @@ void HipMemcpyWithStreamMultiThreadtests::
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
HIP_CHECK(hipGetLastError());
}
for (int i=0; i < numDevices; ++i) {
projects/hip-tests/catch/unit/memory/hipMemoryAllocateCoherent.cc
+1
View File
@@ -53,6 +53,7 @@ TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
std::cout << clkRate << std::endl;
hipLaunchKernelGGL(HIP_KERNEL_NAME(Kernel), dim3(1), dim3(blocks),
0, 0, hostRes, clkRate);
HIP_CHECK(hipGetLastError());
int eleCounter = 0;
while (eleCounter < blocks) {
// blocks until the value changes
projects/hip-tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc
+2 -1
View File
@@ -31,7 +31,7 @@
/* Defines */
#define NUM_THREADS 1000
#define ITER 100
#define ITER 10
#define NUM_H 256
#define NUM_W 256
@@ -105,6 +105,7 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") {
for (size_t k = 0; k < ITER; k++) {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, hipStreamPerThread, B_d, C_d, elements);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H,
hipStreamPerThread));
projects/hip-tests/catch/unit/memory/hipMemset3DFunctional.cc
+1 -1
View File
@@ -369,7 +369,7 @@ static void seekAndSet3DArrayPortion(bool bAsync) {
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
arr_dimensions.x, arr_dimensions.y);
arr_dimensions.y, arr_dimensions.z);
myparms.srcPtr = devicePitchedPointer;
myparms.extent = extent;
#if HT_NVIDIA
projects/hip-tests/catch/unit/memory/hipMemset3DRegressMultiThread.cc
+1
View File
@@ -84,6 +84,7 @@ static void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr,
hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0,
stream, devpPtr, extent, memsetval);
HIP_CHECK(hipGetLastError());
HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream));
HIPCHECK(hipMemcpy3DAsync(&myparms, stream));
}
projects/hip-tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc
+4
View File
@@ -92,6 +92,7 @@ static bool testhipMemsetAsyncWithKernel(bool UseStrmPerThrd) {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, hipStreamPerThread, obj.B_d,
obj.C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal, N, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
@@ -117,6 +118,7 @@ static bool testhipMemsetD32AsyncWithKernel() {
for (int k = 0 ; k < ITER ; k++) {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
N, obj.stream));
HIP_CHECK(hipStreamSynchronize(obj.stream));
@@ -135,6 +137,7 @@ static bool testhipMemsetD16AsyncWithKernel() {
for (int k = 0 ; k < ITER ; k++) {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)obj.C_d , obj.memSetVal,
N, obj.stream));
HIP_CHECK(hipStreamSynchronize(obj.stream));
@@ -153,6 +156,7 @@ static bool testhipMemsetD8AsyncWithKernel() {
for (int k = 0; k < ITER; k++) {
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)obj.C_d, obj.memSetVal,
N, obj.stream));
HIP_CHECK(hipStreamSynchronize(obj.stream));
projects/hip-tests/catch/unit/memory/hipPointerGetAttribute.cc
+1
View File
@@ -138,6 +138,7 @@ TEST_CASE("Unit_hipPointerGetAttribute_KernelUpdation") {
reinterpret_cast<hipDeviceptr_t>(A_d)));
hipLaunchKernelGGL(var_update, dim3(1), dim3(1), 0, 0,
reinterpret_cast<int *>(data));
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
for (unsigned int i = 0; i < N; i++) {