From ac4267e21e300c2a01b7db84bcb223860818e6ac Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 14 Mar 2023 21:42:20 +0530 Subject: [PATCH] SWDEV-1 - Fix warnings while building tests (#201) Change-Id: I7faf3446f8c9dafbbfc6bab4c6c1b756f0e4219d --- catch/include/utils.hh | 2 +- catch/unit/callback/hipGetStreamDeviceId.cc | 2 +- catch/unit/deviceLib/SimpleAtomicsTest.cc | 14 ++++++------- catch/unit/deviceLib/hipStdComplex.cc | 8 +++++--- catch/unit/deviceLib/hipTestAtomicAdd.cc | 4 +--- catch/unit/graph/hipLaunchHostFunc.cc | 4 ++-- catch/unit/graph/hipStreamBeginCapture.cc | 20 +++++++++---------- catch/unit/graph/hipStreamEndCapture.cc | 2 +- catch/unit/graph/hipStreamGetCaptureInfo.cc | 6 ++++-- .../unit/graph/hipStreamGetCaptureInfo_v2.cc | 8 +++++--- catch/unit/graph/hipStreamIsCapturing.cc | 4 ++-- .../hipStreamUpdateCaptureDependencies.cc | 8 ++++---- catch/unit/graph/stream_capture_common.hh | 3 +++ 13 files changed, 46 insertions(+), 39 deletions(-) diff --git a/catch/include/utils.hh b/catch/include/utils.hh index bb7fce06e4..fec888bd97 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -40,7 +40,7 @@ void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements template void ArrayFindIfNot(It begin, It end, const T expected_value) { const auto it = std::find_if_not( - begin, end, [expected_value](const int elem) { return expected_value == elem; }); + begin, end, [expected_value](const T elem) { return expected_value == elem; }); if (it != end) { const auto idx = std::distance(begin, it); diff --git a/catch/unit/callback/hipGetStreamDeviceId.cc b/catch/unit/callback/hipGetStreamDeviceId.cc index e8d718387e..dc2a36a08d 100644 --- a/catch/unit/callback/hipGetStreamDeviceId.cc +++ b/catch/unit/callback/hipGetStreamDeviceId.cc @@ -70,7 +70,7 @@ TEST_CASE("Unit_hipGetStreamDeviceId_Positive_Multithreaded_Basic") { const int device_count = HipTest::getDeviceCount(); auto thread_function = [&]() { - for(unsigned int id = 0; id < device_count; ++id) { + for(int id = 0; id < device_count; ++id) { HIP_CHECK_THREAD(hipSetDevice(id)); StreamGuard stream_guard{Streams::perThread}; diff --git a/catch/unit/deviceLib/SimpleAtomicsTest.cc b/catch/unit/deviceLib/SimpleAtomicsTest.cc index db481e1e22..927344af36 100644 --- a/catch/unit/deviceLib/SimpleAtomicsTest.cc +++ b/catch/unit/deviceLib/SimpleAtomicsTest.cc @@ -108,9 +108,9 @@ template{}>::type* = nullptr> bool VerifyIntegral(T* gpuData, int len) { // atomic Max T val = 0; - for (T i = 0; i < len; ++i) { + for (int i = 0; i < len; ++i) { // fourth element should be len-1 - val = max(val, i); + val = max(val, static_cast(i)); } REQUIRE(val == gpuData[3]); @@ -118,14 +118,14 @@ bool VerifyIntegral(T* gpuData, int len) { // atomic Min val = 1 << 8; - for (T i = 0; i < len; ++i) { - val = min(val, i); + for (int i = 0; i < len; ++i) { + val = min(val, static_cast(i)); } REQUIRE(val == gpuData[4]); // atomic Inc - int limit = 17; + T limit = 17; val = 0; for (int i = 0; i < len; ++i) { @@ -145,9 +145,9 @@ bool VerifyIntegral(T* gpuData, int len) { REQUIRE(val == gpuData[6]); // atomic CAS - for (T i = 0; i < len; ++i) { + for (int i = 0; i < len; ++i) { // eighth element should be a member of [0, len) - if (i == gpuData[7]) { + if (static_cast(i) == gpuData[7]) { return true; break; } diff --git a/catch/unit/deviceLib/hipStdComplex.cc b/catch/unit/deviceLib/hipStdComplex.cc index 8cde8de8ed..f900625d89 100644 --- a/catch/unit/deviceLib/hipStdComplex.cc +++ b/catch/unit/deviceLib/hipStdComplex.cc @@ -24,7 +24,6 @@ THE SOFTWARE. // Tolerance for error const double tolerance = 1e-6; -const bool verbose = false; #define LEN 64 @@ -117,14 +116,15 @@ void test() { hipLaunchKernelGGL(kernel, dim3(1), dim3(LEN), 0, 0, Ad, Bd, Cd, CK); HIP_CHECK(hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost)); + bool pass = true; for (int i = 0; i < LEN; i++) { ComplexT Expected = calc(A[i], B[i], CK); FloatT error = abs(C[i] - Expected); if (abs(Expected) > tolerance) error /= abs(Expected); - bool pass = error < tolerance; + pass &= error < tolerance; } - return true; + return pass; }; #define OP(x) assert(test_fun(CK_##x)); @@ -140,6 +140,7 @@ void test() { delete[] D; } +#if HT_AMD TEST_CASE("Unit_StdComplex") { SECTION("Test run with float") { test(); @@ -148,3 +149,4 @@ TEST_CASE("Unit_StdComplex") { test(); } } +#endif diff --git a/catch/unit/deviceLib/hipTestAtomicAdd.cc b/catch/unit/deviceLib/hipTestAtomicAdd.cc index 4834965942..c728d85b18 100644 --- a/catch/unit/deviceLib/hipTestAtomicAdd.cc +++ b/catch/unit/deviceLib/hipTestAtomicAdd.cc @@ -56,13 +56,11 @@ bool p_atomicNoRet = false; template __global__ void atomicnoret_manywaves(T* C_d) { - size_t tid = (blockIdx.x * blockDim.x + threadIdx.x); atomicAddNoRet(C_d, INCREMENT_VALUE); } template __global__ void atomic_manywaves(T* C_d) { - size_t tid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); atomicAdd(C_d, INCREMENT_VALUE); } @@ -95,7 +93,7 @@ bool atomictest_manywaves(const T& initial_val) { // Copy result from device to host HIP_CHECK(hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost)); REQUIRE(hOData[0] == initial_val+ - (INCREMENT_VALUE*(ThreadsperBlock*numBlocks))); + static_cast(INCREMENT_VALUE*(ThreadsperBlock*numBlocks))); // Cleanup memory free(hOData); diff --git a/catch/unit/graph/hipLaunchHostFunc.cc b/catch/unit/graph/hipLaunchHostFunc.cc index 005bb6b736..1d9ea95b00 100644 --- a/catch/unit/graph/hipLaunchHostFunc.cc +++ b/catch/unit/graph/hipLaunchHostFunc.cc @@ -117,7 +117,7 @@ TEST_CASE("Unit_hipLaunchHostFunc_Positive_Functional") { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), 1, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -171,7 +171,7 @@ TEST_CASE("Unit_hipLaunchHostFunc_Positive_Thread") { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), 1, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); diff --git a/catch/unit/graph/hipStreamBeginCapture.cc b/catch/unit/graph/hipStreamBeginCapture.cc index e7e8b72761..81f97e5f92 100644 --- a/catch/unit/graph/hipStreamBeginCapture.cc +++ b/catch/unit/graph/hipStreamBeginCapture.cc @@ -84,7 +84,7 @@ void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStre REQUIRE(graphExec != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -243,7 +243,7 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea REQUIRE(graphExec2 != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { // Execute the Graphs HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); @@ -286,7 +286,7 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t& REQUIRE(graphExec2 != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { // Execute the Graphs HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); @@ -340,15 +340,15 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea REQUIRE(graphExec2 != nullptr); // Execute the Graphs - for (int iter = 0; iter < kLaunchIters; iter++) { + for (size_t iter = 0; iter < kLaunchIters; iter++) { std::fill_n(A_h.host_ptr(), N, iter); std::fill_n(C_h.host_ptr(), N, iter); HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); HIP_CHECK(hipStreamSynchronize(stream1)); HIP_CHECK(hipStreamSynchronize(stream2)); - ArrayFindIfNot(B_h.host_ptr(), iter * iter, N); - ArrayFindIfNot(D_h.host_ptr(), iter * iter, N); + ArrayFindIfNot(B_h.host_ptr(), static_cast(iter * iter), N); + ArrayFindIfNot(D_h.host_ptr(), static_cast(iter * iter), N); } // Free @@ -407,15 +407,15 @@ static void multithreadedTest(hipStreamCaptureMode mode) { REQUIRE(graphExec2 != nullptr); // Execute the Graphs - for (int iter = 0; iter < kLaunchIters; iter++) { + for (size_t iter = 0; iter < kLaunchIters; iter++) { std::fill_n(A_h.host_ptr(), N, iter); std::fill_n(C_h.host_ptr(), N, iter); HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); HIP_CHECK(hipStreamSynchronize(stream1)); HIP_CHECK(hipStreamSynchronize(stream2)); - ArrayFindIfNot(B_h.host_ptr(), iter * iter, N); - ArrayFindIfNot(D_h.host_ptr(), iter * iter, N); + ArrayFindIfNot(B_h.host_ptr(), static_cast(iter * iter), N); + ArrayFindIfNot(D_h.host_ptr(), static_cast(iter * iter), N); } // Free @@ -1212,7 +1212,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") { hipGraphExec_t graphExec{nullptr}; HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Verify graph - for (int iter = 0; iter < kLaunchIters; iter++) { + for (size_t iter = 0; iter < kLaunchIters; iter++) { std::fill_n(Ah.host_ptr(), N, iter); std::fill_n(Bh.host_ptr(), N, iter); HIP_CHECK(hipGraphLaunch(graphExec, streams[0])); diff --git a/catch/unit/graph/hipStreamEndCapture.cc b/catch/unit/graph/hipStreamEndCapture.cc index 22762268ab..39f25ce2a6 100644 --- a/catch/unit/graph/hipStreamEndCapture.cc +++ b/catch/unit/graph/hipStreamEndCapture.cc @@ -194,7 +194,7 @@ TEST_CASE("Unit_hipStreamEndCapture_Positive_Thread") { HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); diff --git a/catch/unit/graph/hipStreamGetCaptureInfo.cc b/catch/unit/graph/hipStreamGetCaptureInfo.cc index 77c3dac362..d8f8cb5d55 100644 --- a/catch/unit/graph/hipStreamGetCaptureInfo.cc +++ b/catch/unit/graph/hipStreamGetCaptureInfo.cc @@ -71,7 +71,7 @@ void checkStreamCaptureInfo(hipStreamCaptureMode mode, hipStream_t stream) { REQUIRE(graphExec != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -166,7 +166,9 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Positive_UniqueID") { * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamGetCaptureInfo_Negative_Parameters") { +#if HT_NVIDIA hipStreamCaptureStatus cStatus; +#endif unsigned long long capSequenceID; // NOLINT const auto stream_type = GENERATE(Streams::perThread, Streams::created); StreamGuard stream_guard(stream_type); @@ -194,4 +196,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Negative_Parameters") { hipErrorContextIsDestroyed); } #endif -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc b/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc index 5eddd04655..ea67318ef9 100644 --- a/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc +++ b/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc @@ -40,7 +40,7 @@ void checkStreamCaptureInfo_v2(hipStreamCaptureMode mode, hipStream_t stream) { hipGraph_t graph{nullptr}, capInfoGraph{nullptr}; hipGraphExec_t graphExec{nullptr}; const hipGraphNode_t* nodelist{}; - int numDepsCreated = 0; + size_t numDepsCreated = 0; hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; hipGraphNodeType type(hipGraphNodeTypeEmpty); unsigned long long capSequenceID = 0; // NOLINT @@ -119,7 +119,7 @@ void checkStreamCaptureInfo_v2(hipStreamCaptureMode mode, hipStream_t stream) { REQUIRE(graphExec != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -221,7 +221,9 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Positive_UniqueID") { */ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Negative_Parameters") { hipGraph_t capInfoGraph{}; +#if HT_NVIDIA hipStreamCaptureStatus captureStatus; +#endif unsigned long long capSequenceID; // NOLINT size_t numDependencies; const hipGraphNode_t* nodelist{}; @@ -256,4 +258,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Negative_Parameters") { hipErrorContextIsDestroyed); } #endif -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipStreamIsCapturing.cc b/catch/unit/graph/hipStreamIsCapturing.cc index 1c681b4243..c6a77c316e 100644 --- a/catch/unit/graph/hipStreamIsCapturing.cc +++ b/catch/unit/graph/hipStreamIsCapturing.cc @@ -133,7 +133,7 @@ void checkStreamCaptureStatus(hipStreamCaptureMode mode, hipStream_t stream) { REQUIRE(graphExec != nullptr); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); @@ -209,4 +209,4 @@ TEST_CASE("Unit_hipStreamIsCapturing_Positive_Thread") { HIP_CHECK(hipStreamEndCapture(stream, &graph)); HIP_CHECK(hipGraphDestroy(graph)); -} \ No newline at end of file +} diff --git a/catch/unit/graph/hipStreamUpdateCaptureDependencies.cc b/catch/unit/graph/hipStreamUpdateCaptureDependencies.cc index 871ede9ea4..e35dd317d6 100644 --- a/catch/unit/graph/hipStreamUpdateCaptureDependencies.cc +++ b/catch/unit/graph/hipStreamUpdateCaptureDependencies.cc @@ -34,7 +34,7 @@ THE SOFTWARE. * update the set of dependencies in a capturing stream */ -static __global__ void vectorSet(const float* A_d, float* B_d, int64_t NELEM) { +static __global__ void vectorSet(const float* A_d, float* B_d, size_t NELEM) { size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; @@ -159,7 +159,7 @@ static void UpdateStreamCaptureDependenciesSet(hipStream_t stream, HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); std::fill_n(C_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); @@ -275,7 +275,7 @@ static void UpdateStreamCaptureDependenciesAdd(hipStream_t stream, HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); // Replay the recorded sequence multiple times - for (int i = 0; i < kLaunchIters; i++) { + for (size_t i = 0; i < kLaunchIters; i++) { std::fill_n(A_h.host_ptr(), N, static_cast(i)); std::fill_n(C_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); @@ -367,7 +367,7 @@ TEST_CASE("Unit_hipStreamUpdateCaptureDependencies_Positive_Parameters") { const hipStreamUpdateCaptureDependenciesFlags flag = GENERATE(hipStreamAddCaptureDependencies, hipStreamSetCaptureDependencies); - HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamBeginCapture(stream, captureMode)); //hipStreamCaptureModeGlobal)); HIP_CHECK(hipStreamUpdateCaptureDependencies(stream, nullptr, 0, flag)); diff --git a/catch/unit/graph/stream_capture_common.hh b/catch/unit/graph/stream_capture_common.hh index 2e1fe9bfaa..71a6548c17 100644 --- a/catch/unit/graph/stream_capture_common.hh +++ b/catch/unit/graph/stream_capture_common.hh @@ -40,6 +40,8 @@ void captureSequenceSimple(T* hostMem1, T* devMem1, T* hostMem2, size_t N, template void captureSequenceLinear(T* hostMem1, T* devMem1, T* hostMem2, T* devMem2, size_t N, hipStream_t captureStream) { + + {(void)(hostMem2);} //unused hostMem2 size_t Nbytes = N * sizeof(T); HIP_CHECK(hipMemcpyAsync(devMem1, hostMem1, Nbytes, hipMemcpyHostToDevice, captureStream)); @@ -51,6 +53,7 @@ template void captureSequenceBranched(T* hostMem1, T* devMem1, T* hostMem2, T* devMem2, size_t N, hipStream_t captureStream, std::vector& streams, std::vector& events) { + {(void)(hostMem2);} //unused hostMem2 size_t Nbytes = N * sizeof(T); HIP_CHECK(hipEventRecord(events[0], captureStream));