From cea96af7597ebd3e0774ecad7290b432eadffc1e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 26 Oct 2022 03:59:02 +0000 Subject: [PATCH 01/14] SWDEV-355313 - Move catch tests and samples Change-Id: I66f0c09e9c7405ec7430b1883e0e89542fdb87a0 --- LICENSE.txt | 20 + catch/ABM/AddKernels/CMakeLists.txt | 8 + catch/ABM/AddKernels/add.cc | 41 + catch/ABM/CMakeLists.txt | 1 + catch/CMakeLists.txt | 249 + catch/README.md | 204 + catch/TypeQualifiers/CMakeLists.txt | 8 + catch/TypeQualifiers/hipManagedKeyword.cc | 76 + catch/catchProp_in_rc.in | 40 + catch/external/Catch2/LICENSE.txt | 23 + catch/external/Catch2/catch.hpp | 17881 ++++++++++++++++ .../external/Catch2/cmake/Catch2/Catch.cmake | 249 + .../Catch2/cmake/Catch2/Catch2Config.cmake | 34 + .../cmake/Catch2/Catch2ConfigVersion.cmake | 51 + .../Catch2/cmake/Catch2/Catch2Targets.cmake | 99 + .../Catch2/cmake/Catch2/CatchAddTests.cmake | 138 + .../cmake/Catch2/ParseAndAddCatchTests.cmake | 252 + .../cmake/Catch2/catch_include.cmake.in | 34 + catch/external/picojson/LICENSE | 25 + catch/external/picojson/picojson.h | 1200 ++ catch/hipTestMain/CMakeLists.txt | 30 + .../config/config_amd_linux_MI2xx.json | 8 + .../config/config_amd_linux_common.json | 7 + .../config/config_amd_windows_MI2xx.json | 76 + .../config/config_amd_windows_common.json | 83 + catch/hipTestMain/hip_test_context.cc | 272 + catch/hipTestMain/main.cc | 16 + catch/hipTestMain/standalone_main.cc | 2 + catch/include/hip_test_checkers.hh | 398 + catch/include/hip_test_common.hh | 411 + catch/include/hip_test_context.hh | 173 + catch/include/hip_test_filesystem.hh | 89 + catch/include/hip_test_helper.hh | 70 + catch/include/hip_test_kernels.hh | 107 + catch/include/hip_test_process.hh | 117 + catch/include/hip_test_rtc.hh | 283 + catch/include/hip_test_smi.hh | 88 + catch/include/hip_texture_helper.hh | 370 + catch/include/kernel_mapping.hh | 27 + catch/include/kernels.hh | 55 + catch/include/resource_guards.hh | 125 + catch/include/threaded_zig_zag_test.hh | 110 + catch/include/utils.hh | 102 + catch/kernels/CMakeLists.txt | 8 + catch/kernels/Set.cpp | 6 + catch/kernels/vectorADD.inl | 10 + catch/multiproc/CMakeLists.txt | 35 + catch/multiproc/childMalloc.cc | 62 + catch/multiproc/deviceAllocationMproc.cc | 319 + catch/multiproc/dummy_kernel.cpp | 26 + .../hipDeviceComputeCapabilityMproc.cc | 159 + catch/multiproc/hipDeviceGetPCIBusIdMproc.cc | 258 + catch/multiproc/hipDeviceTotalMemMproc.cc | 161 + catch/multiproc/hipGetDeviceAttributeMproc.cc | 164 + catch/multiproc/hipGetDeviceCountMproc.cc | 54 + .../multiproc/hipGetDevicePropertiesMproc.cc | 165 + catch/multiproc/hipIpcEventHandle.cc | 359 + catch/multiproc/hipIpcMemAccessTest.cc | 222 + catch/multiproc/hipMallocConcurrencyMproc.cc | 244 + catch/multiproc/hipMemCoherencyTstMProc.cc | 781 + catch/multiproc/hipNoGpuTsts.cc | 1687 ++ catch/multiproc/hipSetGetDeviceMproc.cc | 570 + catch/packaging/hip-tests.txt | 107 + catch/stress/CMakeLists.txt | 9 + catch/stress/deviceallocation/CMakeLists.txt | 8 + .../Stress_deviceAllocationStress.cc | 487 + catch/stress/memory/CMakeLists.txt | 12 + catch/stress/memory/hipHostMalloc.cc | 52 + catch/stress/memory/hipMalloc.cc | 50 + catch/stress/memory/hipMallocManagedStress.cc | 330 + .../memory/hipMemPrftchAsyncStressTst.cc | 133 + catch/stress/memory/hipMemcpyMThreadMSize.cc | 275 + catch/stress/memory/memcpy.cc | 34 + catch/stress/printf/CMakeLists.txt | 9 + .../printf/Stress_printf_ComplexKernels.cc | 517 + .../printf/Stress_printf_SimpleKernels.cc | 795 + catch/stress/printf/printf_common.h | 99 + catch/stress/stream/CMakeLists.txt | 10 + catch/stress/stream/Stress_hipStreamCreate.cc | 203 + catch/stress/stream/streamEnqueue.cc | 238 + catch/unit/CMakeLists.txt | 36 + catch/unit/clock/CMakeLists.txt | 29 + catch/unit/clock/hipClockCheck.cc | 116 + .../coalesced_groups_shfl_down.cc | 267 + .../coalesced_groups_shfl_up.cc | 251 + .../simple_coalesced_groups.cc | 576 + catch/unit/device/CMakeLists.txt | 36 + catch/unit/device/getDeviceCount_exe.cc | 89 + catch/unit/device/hipChooseDevice.cc | 57 + .../unit/device/hipDeviceComputeCapability.cc | 71 + catch/unit/device/hipDeviceGetByPCIBusId.cc | 130 + catch/unit/device/hipDeviceGetLimit.cc | 51 + catch/unit/device/hipDeviceGetName.cc | 137 + catch/unit/device/hipDeviceGetP2PAttribute.cc | 126 + .../device/hipDeviceGetP2PAttribute_exe.cc | 86 + catch/unit/device/hipDeviceGetPCIBusId.cc | 117 + catch/unit/device/hipDeviceGetUuid.cc | 51 + .../unit/device/hipDeviceSetGetCacheConfig.cc | 45 + catch/unit/device/hipDeviceSetLimit.cc | 101 + catch/unit/device/hipDeviceSynchronize.cc | 79 + catch/unit/device/hipDeviceTotalMem.cc | 99 + .../device/hipExtGetLinkTypeAndHopCount.cc | 100 + catch/unit/device/hipGetDeviceAttribute.cc | 558 + catch/unit/device/hipGetDeviceCount.cc | 64 + catch/unit/device/hipGetDeviceProperties.cc | 321 + catch/unit/device/hipGetSetDeviceFlags.cc | 147 + catch/unit/device/hipRuntimeGetVersion.cc | 34 + catch/unit/device/hipSetGetDevice.cc | 155 + .../AtomicAdd_Coherent_withnoUnsafeflag.cc | 92 + .../AtomicAdd_Coherent_withoutflag.cc | 91 + .../AtomicAdd_Coherent_withunsafeflag.cc | 100 + .../AtomicAdd_NonCoherent_withnoUnsafeflag.cc | 92 + .../AtomicAdd_NonCoherent_withoutflag.cc | 92 + .../AtomicAdd_NonCoherent_withunsafeflag.cc | 99 + catch/unit/deviceLib/BuiltIns_fadd.cc | 294 + catch/unit/deviceLib/BuiltIns_fmax.cc | 357 + catch/unit/deviceLib/BuiltIns_fmin.cc | 360 + catch/unit/deviceLib/CMakeLists.txt | 98 + catch/unit/deviceLib/anyAll.cc | 84 + catch/unit/deviceLib/ballot.cc | 86 + catch/unit/deviceLib/bitExtract.cc | 195 + catch/unit/deviceLib/bitInsert.cc | 217 + catch/unit/deviceLib/brev.cc | 151 + catch/unit/deviceLib/clz.cc | 173 + catch/unit/deviceLib/defs.h | 36 + catch/unit/deviceLib/deviceAllocCommon.h | 132 + catch/unit/deviceLib/deviceAllocation.cc | 1494 ++ catch/unit/deviceLib/ffs.cc | 147 + catch/unit/deviceLib/floatMath.cc | 66 + catch/unit/deviceLib/floatTM.cc | 227 + catch/unit/deviceLib/funnelshift.cc | 208 + catch/unit/deviceLib/hipTestDeviceSymbol.cc | 228 + catch/unit/deviceLib/kerDevAllocMultCO.cc | 39 + catch/unit/deviceLib/kerDevAllocSingleKer.cc | 57 + catch/unit/deviceLib/kerDevFreeMultCO.cc | 47 + catch/unit/deviceLib/kerDevWriteMultCO.cc | 36 + catch/unit/deviceLib/ldg.cc | 251 + catch/unit/deviceLib/mbcnt.cc | 103 + catch/unit/deviceLib/popc.cc | 136 + catch/unit/deviceLib/syncthreadsand.cc | 127 + catch/unit/deviceLib/syncthreadscount.cc | 138 + catch/unit/deviceLib/syncthreadsor.cc | 128 + catch/unit/deviceLib/threadfence_system.cc | 113 + catch/unit/deviceLib/unsafeAtomicAdd.cc | 133 + ...safeAtomicAdd_Coherent_withnounsafeflag.cc | 100 + .../unsafeAtomicAdd_Coherent_withoutflag.cc | 100 + ...unsafeAtomicAdd_Coherent_withunsafeflag.cc | 101 + ...eAtomicAdd_NonCoherent_withnounsafeflag.cc | 98 + ...unsafeAtomicAdd_NonCoherent_withoutflag.cc | 98 + ...afeAtomicAdd_NonCoherent_withunsafeflag.cc | 98 + catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc | 583 + catch/unit/deviceLib/vectorTypesDevice.cc | 271 + catch/unit/event/CMakeLists.txt | 23 + catch/unit/event/Unit_hipEvent.cc | 189 + catch/unit/event/Unit_hipEventElapsedTime.cc | 132 + catch/unit/event/Unit_hipEventIpc.cc | 104 + catch/unit/event/Unit_hipEventQuery.cc | 97 + catch/unit/event/Unit_hipEventRecord.cc | 141 + catch/unit/event/Unit_hipEvent_Negative.cc | 82 + catch/unit/event/hipEventCreateWithFlags.cc | 44 + catch/unit/event/hipEventDestroy.cc | 115 + catch/unit/event/hipEventSynchronize.cc | 130 + catch/unit/graph/CMakeLists.txt | 85 + catch/unit/graph/hipGraph.cc | 346 + catch/unit/graph/hipGraphAddChildGraphNode.cc | 424 + catch/unit/graph/hipGraphAddDependencies.cc | 249 + catch/unit/graph/hipGraphAddEmptyNode.cc | 100 + .../unit/graph/hipGraphAddEventRecordNode.cc | 346 + catch/unit/graph/hipGraphAddHostNode.cc | 314 + catch/unit/graph/hipGraphAddKernelNode.cc | 101 + catch/unit/graph/hipGraphAddMemcpyNode.cc | 520 + catch/unit/graph/hipGraphAddMemcpyNode1D.cc | 200 + .../graph/hipGraphAddMemcpyNodeFromSymbol.cc | 437 + .../graph/hipGraphAddMemcpyNodeToSymbol.cc | 402 + catch/unit/graph/hipGraphAddMemsetNode.cc | 123 + .../graph/hipGraphChildGraphNodeGetGraph.cc | 163 + catch/unit/graph/hipGraphClone.cc | 320 + catch/unit/graph/hipGraphDestroyNode.cc | 139 + .../graph/hipGraphEventRecordNodeGetEvent.cc | 130 + .../graph/hipGraphEventRecordNodeSetEvent.cc | 246 + .../graph/hipGraphEventWaitNodeGetEvent.cc | 130 + .../graph/hipGraphEventWaitNodeSetEvent.cc | 317 + .../hipGraphExecChildGraphNodeSetParams.cc | 347 + .../hipGraphExecEventRecordNodeSetEvent.cc | 289 + .../hipGraphExecEventWaitNodeSetEvent.cc | 313 + .../graph/hipGraphExecHostNodeSetParams.cc | 276 + .../graph/hipGraphExecKernelNodeSetParams.cc | 186 + .../graph/hipGraphExecMemcpyNodeSetParams.cc | 259 + .../hipGraphExecMemcpyNodeSetParams1D.cc | 198 + ...pGraphExecMemcpyNodeSetParamsFromSymbol.cc | 303 + ...hipGraphExecMemcpyNodeSetParamsToSymbol.cc | 316 + .../graph/hipGraphExecMemsetNodeSetParams.cc | 201 + catch/unit/graph/hipGraphExecUpdate.cc | 341 + catch/unit/graph/hipGraphGetEdges.cc | 251 + catch/unit/graph/hipGraphGetNodes.cc | 219 + catch/unit/graph/hipGraphGetRootNodes.cc | 239 + catch/unit/graph/hipGraphHostNodeGetParams.cc | 275 + catch/unit/graph/hipGraphHostNodeSetParams.cc | 232 + catch/unit/graph/hipGraphInstantiate.cc | 75 + .../graph/hipGraphInstantiateWithFlags.cc | 313 + .../unit/graph/hipGraphKernelNodeGetParams.cc | 159 + .../unit/graph/hipGraphKernelNodeSetParams.cc | 145 + catch/unit/graph/hipGraphLaunch.cc | 307 + .../unit/graph/hipGraphMemcpyNodeGetParams.cc | 232 + .../unit/graph/hipGraphMemcpyNodeSetParams.cc | 215 + .../graph/hipGraphMemcpyNodeSetParams1D.cc | 185 + .../hipGraphMemcpyNodeSetParamsFromSymbol.cc | 260 + .../hipGraphMemcpyNodeSetParamsToSymbol.cc | 264 + .../unit/graph/hipGraphMemsetNodeGetParams.cc | 141 + .../unit/graph/hipGraphMemsetNodeSetParams.cc | 185 + catch/unit/graph/hipGraphNodeFindInClone.cc | 241 + .../unit/graph/hipGraphNodeGetDependencies.cc | 368 + .../graph/hipGraphNodeGetDependentNodes.cc | 366 + catch/unit/graph/hipGraphNodeGetType.cc | 194 + .../unit/graph/hipGraphRemoveDependencies.cc | 462 + catch/unit/graph/hipSimpleGraphWithKernel.cc | 160 + catch/unit/graph/hipStreamBeginCapture.cc | 228 + catch/unit/graph/hipStreamEndCapture.cc | 176 + catch/unit/graph/hipStreamGetCaptureInfo.cc | 231 + catch/unit/graph/hipStreamIsCapturing.cc | 215 + catch/unit/kernel/CMakeLists.txt | 29 + catch/unit/kernel/hipLaunchBounds.cc | 173 + .../unit/kernel/hipMemFaultStackAllocation.cc | 102 + catch/unit/memory/CMakeLists.txt | 186 + catch/unit/memory/DriverContext.cc | 36 + catch/unit/memory/DriverContext.hh | 41 + catch/unit/memory/MemUtils.hh | 407 + catch/unit/memory/hipArray.cc | 76 + catch/unit/memory/hipArray3DCreate.cc | 345 + catch/unit/memory/hipArrayCommon.hh | 221 + catch/unit/memory/hipArrayCreate.cc | 363 + catch/unit/memory/hipDrvMemcpy3D.cc | 573 + catch/unit/memory/hipDrvMemcpy3DAsync.cc | 594 + catch/unit/memory/hipDrvPtrGetAttributes.cc | 176 + catch/unit/memory/hipFree.cc | 421 + catch/unit/memory/hipHmmOvrSubscriptionTst.cc | 213 + catch/unit/memory/hipHostGetDevicePointer.cc | 78 + catch/unit/memory/hipHostGetFlags.cc | 223 + catch/unit/memory/hipHostMalloc.cc | 244 + catch/unit/memory/hipHostMallocTests.cc | 68 + catch/unit/memory/hipHostRegister.cc | 224 + catch/unit/memory/hipHostUnregister.cc | 91 + catch/unit/memory/hipMalloc3D.cc | 113 + catch/unit/memory/hipMalloc3DArray.cc | 435 + catch/unit/memory/hipMallocArray.cc | 710 + catch/unit/memory/hipMallocConcurrency.cc | 426 + catch/unit/memory/hipMallocManaged.cc | 140 + catch/unit/memory/hipMallocManagedCommon.hh | 26 + catch/unit/memory/hipMallocManagedFlagsTst.cc | 229 + .../memory/hipMallocManaged_MultiScenario.cc | 406 + catch/unit/memory/hipMallocMngdMultiThread.cc | 507 + catch/unit/memory/hipMallocPitch.cc | 556 + catch/unit/memory/hipMemAdvise.cc | 927 + catch/unit/memory/hipMemAdviseMmap.cc | 87 + catch/unit/memory/hipMemCoherencyTst.cc | 258 + catch/unit/memory/hipMemGetInfo.cc | 576 + catch/unit/memory/hipMemPoolApi.cc | 554 + catch/unit/memory/hipMemPrefetchAsync.cc | 124 + .../unit/memory/hipMemPrefetchAsyncExtTsts.cc | 348 + catch/unit/memory/hipMemPtrGetInfo.cc | 54 + catch/unit/memory/hipMemRangeGetAttribute.cc | 408 + catch/unit/memory/hipMemVmm.cc | 85 + catch/unit/memory/hipMemcpy.cc | 619 + catch/unit/memory/hipMemcpy2D.cc | 284 + catch/unit/memory/hipMemcpy2DAsync.cc | 409 + catch/unit/memory/hipMemcpy2DFromArray.cc | 331 + .../unit/memory/hipMemcpy2DFromArrayAsync.cc | 372 + catch/unit/memory/hipMemcpy2DToArray.cc | 332 + catch/unit/memory/hipMemcpy2DToArrayAsync.cc | 369 + catch/unit/memory/hipMemcpy3D.cc | 627 + catch/unit/memory/hipMemcpy3DAsync.cc | 755 + catch/unit/memory/hipMemcpyAllApiNegative.cc | 327 + catch/unit/memory/hipMemcpyAsync.cc | 408 + catch/unit/memory/hipMemcpyAtoH.cc | 219 + catch/unit/memory/hipMemcpyDtoD.cc | 97 + catch/unit/memory/hipMemcpyDtoDAsync.cc | 102 + catch/unit/memory/hipMemcpyFromSymbol.cc | 190 + catch/unit/memory/hipMemcpyHtoA.cc | 229 + catch/unit/memory/hipMemcpyParam2D.cc | 380 + catch/unit/memory/hipMemcpyParam2DAsync.cc | 490 + catch/unit/memory/hipMemcpyPeer.cc | 158 + catch/unit/memory/hipMemcpyPeerAsync.cc | 262 + catch/unit/memory/hipMemcpySync.cc | 227 + catch/unit/memory/hipMemcpyWithStream.cc | 586 + .../memory/hipMemcpyWithStreamMultiThread.cc | 687 + catch/unit/memory/hipMemcpy_MultiThread.cc | 322 + .../unit/memory/hipMemoryAllocateCoherent.cc | 65 + catch/unit/memory/hipMemset.cc | 281 + catch/unit/memory/hipMemset2D.cc | 175 + .../hipMemset2DAsyncMultiThreadAndKernel.cc | 201 + catch/unit/memory/hipMemset3D.cc | 128 + catch/unit/memory/hipMemset3DFunctional.cc | 521 + .../memory/hipMemset3DRegressMultiThread.cc | 267 + catch/unit/memory/hipMemsetAsync.cc | 173 + catch/unit/memory/hipMemsetAsyncAndKernel.cc | 213 + .../unit/memory/hipMemsetAsyncMultiThread.cc | 243 + catch/unit/memory/hipMemsetFunctional.cc | 559 + catch/unit/memory/hipMemsetNegative.cc | 218 + catch/unit/memory/hipMemsetSync.cc | 505 + catch/unit/memory/hipPointerGetAttribute.cc | 357 + catch/unit/memory/hipPointerGetAttributes.cc | 391 + catch/unit/memory/hipPtrGetAttribute.cc | 160 + catch/unit/memory/malloc.cc | 17 + catch/unit/memory/memset.cc | 21 + catch/unit/multiThread/CMakeLists.txt | 10 + .../unit/multiThread/hipMultiThreadDevice.cc | 112 + .../multiThread/hipMultiThreadStreams1.cc | 146 + .../multiThread/hipMultiThreadStreams2.cc | 148 + catch/unit/occupancy/CMakeLists.txt | 9 + ...cupancyMaxActiveBlocksPerMultiprocessor.cc | 91 + .../hipOccupancyMaxPotentialBlockSize.cc | 80 + catch/unit/printf/CMakeLists.txt | 25 + catch/unit/printf/printfFlags.cc | 43 + catch/unit/printf/printfFlags_exe.cc | 42 + catch/unit/printf/printfSpecifiers.cc | 95 + catch/unit/printf/printfSpecifiers_exe.cc | 65 + catch/unit/rtc/CMakeLists.txt | 24 + catch/unit/rtc/customOptions.cc | 181 + catch/unit/rtc/headers/test_header1.h | 4 + catch/unit/rtc/headers/test_header2.h | 4 + catch/unit/rtc/hipRtcFunctional.cc | 87 + catch/unit/rtc/includepath.cc | 128 + catch/unit/rtc/saxpy.cc | 116 + catch/unit/rtc/saxpy.h | 11 + catch/unit/rtc/warpsize.cc | 98 + catch/unit/stream/CMakeLists.txt | 49 + catch/unit/stream/hipAPIStreamDisable.cc | 71 + .../stream/hipDeviceGetStreamPriorityRange.cc | 37 + catch/unit/stream/hipMultiStream.cc | 102 + catch/unit/stream/hipStreamACb_MultiThread.cc | 168 + catch/unit/stream/hipStreamAddCallback.cc | 243 + catch/unit/stream/hipStreamAttachMemAsync.cc | 143 + catch/unit/stream/hipStreamCreate.cc | 35 + catch/unit/stream/hipStreamCreateWithFlags.cc | 85 + .../stream/hipStreamCreateWithPriority.cc | 663 + catch/unit/stream/hipStreamDestroy.cc | 97 + catch/unit/stream/hipStreamGetCUMask.cc | 189 + catch/unit/stream/hipStreamGetFlags.cc | 82 + catch/unit/stream/hipStreamGetPriority.cc | 149 + catch/unit/stream/hipStreamQuery.cc | 141 + catch/unit/stream/hipStreamSynchronize.cc | 169 + catch/unit/stream/hipStreamValue.cc | 559 + catch/unit/stream/hipStreamWaitEvent.cc | 121 + catch/unit/stream/hipStreamWithCUMask.cc | 359 + catch/unit/stream/streamCommon.cc | 101 + catch/unit/stream/streamCommon.hh | 41 + catch/unit/streamperthread/CMakeLists.txt | 12 + .../streamperthread/hipStreamPerThrdTsts.cc | 472 + .../hipStreamPerThread_Basic.cc | 133 + .../hipStreamPerThread_DeviceReset.cc | 99 + .../hipStreamPerThread_Event.cc | 64 + .../hipStreamPerThread_MultiThread.cc | 55 + catch/unit/texture/CMakeLists.txt | 52 + catch/unit/texture/hipBindTex2DPitch.cc | 79 + catch/unit/texture/hipBindTexRef1DFetch.cc | 81 + .../hipCreateTextureObject_ArgValidation.cc | 83 + .../texture/hipCreateTextureObject_Array.cc | 71 + .../texture/hipCreateTextureObject_Linear.cc | 135 + .../texture/hipCreateTextureObject_Pitch2D.cc | 218 + catch/unit/texture/hipGetChanDesc.cc | 45 + .../texture/hipNormalizedFloatValueTex.cc | 161 + .../unit/texture/hipSimpleTexture2DLayered.cc | 109 + catch/unit/texture/hipSimpleTexture3D.cc | 121 + catch/unit/texture/hipTex1DFetchCheckModes.cc | 124 + catch/unit/texture/hipTexObjPitch.cc | 99 + catch/unit/texture/hipTextureMipmapObj2D.cc | 138 + .../unit/texture/hipTextureObj1DCheckModes.cc | 145 + .../texture/hipTextureObj1DCheckSRGBModes.cc | 298 + catch/unit/texture/hipTextureObj1DFetch.cc | 89 + catch/unit/texture/hipTextureObj2D.cc | 107 + .../unit/texture/hipTextureObj2DCheckModes.cc | 157 + .../texture/hipTextureObj2DCheckSRGBModes.cc | 243 + .../unit/texture/hipTextureObj3DCheckModes.cc | 215 + .../unit/texture/hipTextureObjFetchVector.cc | 242 + catch/unit/texture/hipTextureRef2D.cc | 92 + samples/0_Intro/bit_extract/CMakeLists.txt | 52 + samples/0_Intro/bit_extract/Makefile | 46 + samples/0_Intro/bit_extract/README.md | 6 + samples/0_Intro/bit_extract/bit_extract.cpp | 113 + samples/0_Intro/module_api/CMakeLists.txt | 60 + samples/0_Intro/module_api/Makefile | 46 + samples/0_Intro/module_api/defaultDriver.cpp | 87 + .../0_Intro/module_api/launchKernelHcc.cpp | 111 + samples/0_Intro/module_api/runKernel.cpp | 106 + samples/0_Intro/module_api/vcpy_kernel.cpp | 28 + .../0_Intro/module_api_global/CMakeLists.txt | 54 + samples/0_Intro/module_api_global/Makefile | 40 + .../0_Intro/module_api_global/runKernel.cpp | 161 + .../0_Intro/module_api_global/vcpy_kernel.cpp | 38 + samples/0_Intro/square/CMakeLists.txt | 48 + samples/0_Intro/square/Makefile | 47 + samples/0_Intro/square/README.md | 37 + samples/0_Intro/square/square.cu | 98 + samples/0_Intro/square/square.hipref.cpp | 96 + .../1_Utils/hipBusBandwidth/CMakeLists.txt | 44 + samples/1_Utils/hipBusBandwidth/LICENSE.txt | 27 + samples/1_Utils/hipBusBandwidth/Makefile | 43 + .../hipBusBandwidth/ResultDatabase.cpp | 462 + .../1_Utils/hipBusBandwidth/ResultDatabase.h | 89 + .../hipBusBandwidth/hipBusBandwidth.cpp | 1072 + samples/1_Utils/hipCommander/CMakeLists.txt | 55 + samples/1_Utils/hipCommander/LICENSE.txt | 27 + samples/1_Utils/hipCommander/Makefile | 53 + .../1_Utils/hipCommander/ResultDatabase.cpp | 454 + samples/1_Utils/hipCommander/ResultDatabase.h | 89 + samples/1_Utils/hipCommander/c.cmd | 3 + samples/1_Utils/hipCommander/classic.cmd | 1 + samples/1_Utils/hipCommander/hipCommander.cpp | 865 + samples/1_Utils/hipCommander/l2.hcm | 3 + samples/1_Utils/hipCommander/loop.hcm | 3 + samples/1_Utils/hipCommander/loop2.hcm | 2 + .../1_Utils/hipCommander/nullkernel.hip.cpp | 7 + samples/1_Utils/hipCommander/nullkernel.hsaco | Bin 0 -> 39333 bytes .../1_Utils/hipCommander/perf/latency2.hcm | 10 + .../hipCommander/perf/latency_hostsync.hcm | 8 + .../hipCommander/perf/latency_nosync.hcm | 5 + .../hipCommander/perf/latency_nullstream.hcm | 7 + .../perf/modulelaunch_latency.hcm | 5 + .../perf/scripts/latency_2_d2h_h2d.hcm | 9 + .../perf/scripts/latency_2_d2h_kernel.hcm | 9 + .../perf/scripts/latency_2_d2h_sync_h2d.hcm | 9 + .../scripts/latency_2_d2h_sync_kernel.hcm | 9 + .../perf/scripts/latency_2_h2d_d2h.hcm | 9 + .../perf/scripts/latency_2_h2d_kernel.hcm | 9 + .../perf/scripts/latency_2_h2d_sync_d2h.hcm | 9 + .../scripts/latency_2_h2d_sync_kernel.hcm | 9 + .../perf/scripts/latency_2_kernel_d2h.hcm | 9 + .../perf/scripts/latency_2_kernel_h2d.hcm | 9 + .../scripts/latency_2_kernel_sync_d2h.hcm | 9 + .../scripts/latency_2_kernel_sync_h2d.hcm | 9 + .../perf/scripts/latency_2_sync.hcm | 9 + .../perf/scripts/latency_2d2h.hcm | 10 + .../perf/scripts/latency_2d2h_wosync.hcm | 10 + .../perf/scripts/latency_2h2d.hcm | 10 + .../perf/scripts/latency_2h2d_kernel.hcm | 9 + .../scripts/latency_2h2d_kernel_wosync.hcm | 9 + .../perf/scripts/latency_2h2d_wosync.hcm | 10 + .../perf/scripts/latency_2kernels.hcm | 10 + .../perf/scripts/latency_2kernels_wosync.hcm | 10 + .../perf/scripts/latency_3_d2h_h2d.hcm | 9 + .../perf/scripts/latency_3_d2h_kernel.hcm | 9 + .../perf/scripts/latency_3_d2h_sync_h2d.hcm | 9 + .../scripts/latency_3_d2h_sync_kernel.hcm | 9 + .../perf/scripts/latency_3_h2d_d2h.hcm | 10 + .../perf/scripts/latency_3_h2d_kernel.hcm | 9 + .../perf/scripts/latency_3_h2d_sync_d2h.hcm | 9 + .../scripts/latency_3_h2d_sync_kernel.hcm | 9 + .../perf/scripts/latency_3_kernel_d2h.hcm | 9 + .../perf/scripts/latency_3_kernel_h2d.hcm | 9 + .../scripts/latency_3_kernel_sync_d2h.hcm | 9 + .../scripts/latency_3_kernel_sync_h2d.hcm | 9 + .../perf/scripts/latency_3_sync.hcm | 9 + .../perf/scripts/latency_3d2h.hcm | 9 + .../perf/scripts/latency_3d2h_wosync.hcm | 9 + .../perf/scripts/latency_3h2d.hcm | 10 + .../perf/scripts/latency_3h2d_wosync.hcm | 10 + .../perf/scripts/latency_3kernels.hcm | 10 + .../perf/scripts/latency_3kernels_wosync.hcm | 10 + .../perf/scripts/latency_4kernels.hcm | 10 + .../perf/scripts/latency_5kernels.hcm | 10 + .../perf/scripts/latency_6kernels.hcm | 10 + .../hipCommander/perf/scripts/latency_d2h.hcm | 9 + .../perf/scripts/latency_d2h_h2d.hcm | 9 + .../perf/scripts/latency_d2h_kernel.hcm | 9 + .../perf/scripts/latency_d2h_sync_h2d.hcm | 9 + .../perf/scripts/latency_d2h_sync_kernel.hcm | 9 + .../hipCommander/perf/scripts/latency_h2d.hcm | 10 + .../perf/scripts/latency_h2d_10.hcm | 2 + .../perf/scripts/latency_h2d_d2h.hcm | 9 + .../perf/scripts/latency_h2d_kernel.hcm | 9 + .../perf/scripts/latency_h2d_kernel_d2h.hcm | 9 + .../scripts/latency_h2d_kernel_d2h_wosync.hcm | 9 + .../scripts/latency_h2d_kernel_wosync.hcm | 9 + .../perf/scripts/latency_h2d_sync_d2h.hcm | 9 + .../scripts/latency_h2d_sync_kernel_sync.hcm | 9 + .../latency_h2d_sync_kernel_sync_d2h.hcm | 9 + .../perf/scripts/latency_kernel.hcm | 10 + .../perf/scripts/latency_kernel_barrier.hcm | 9 + .../perf/scripts/latency_kernel_d2h.hcm | 9 + .../perf/scripts/latency_kernel_h2d.hcm | 9 + .../perf/scripts/latency_kernel_sync_d2h.hcm | 9 + .../perf/scripts/latency_kernel_sync_h2d.hcm | 9 + .../perf/scripts/latency_streamcreate.hcm | 2 + .../perf/scripts/latency_sync.hcm | 9 + samples/1_Utils/hipCommander/setstream.hcm | 3 + samples/1_Utils/hipCommander/testcase.cpp | 18 + .../1_Utils/hipDispatchLatency/CMakeLists.txt | 59 + samples/1_Utils/hipDispatchLatency/Makefile | 43 + .../hipDispatchEnqueueRateMT.cpp | 196 + .../hipDispatchLatency/hipDispatchLatency.cpp | 141 + .../hipDispatchLatency/test_kernel.cpp | 24 + samples/1_Utils/hipInfo/CMakeLists.txt | 65 + samples/1_Utils/hipInfo/Makefile | 42 + samples/1_Utils/hipInfo/README.md | 6 + samples/1_Utils/hipInfo/hipInfo.cpp | 213 + .../0_MatrixTranspose/CMakeLists.txt | 44 + samples/2_Cookbook/0_MatrixTranspose/Makefile | 59 + .../0_MatrixTranspose/MatrixTranspose.cpp | 122 + .../2_Cookbook/0_MatrixTranspose/Readme.md | 101 + .../2_Cookbook/10_inline_asm/CMakeLists.txt | 44 + samples/2_Cookbook/10_inline_asm/Makefile | 58 + samples/2_Cookbook/10_inline_asm/Readme.md | 60 + .../2_Cookbook/10_inline_asm/inline_asm.cpp | 159 + .../11_texture_driver/CMakeLists.txt | 54 + samples/2_Cookbook/11_texture_driver/Makefile | 40 + .../11_texture_driver/tex2dKernel.cpp | 80 + .../11_texture_driver/texture2dDrv.cpp | 248 + .../CMakeLists.txt | 61 + .../MatrixTranspose.cpp | 122 + .../12_cmake_hip_add_executable/Readme.md | 59 + .../2_Cookbook/13_occupancy/CMakeLists.txt | 44 + samples/2_Cookbook/13_occupancy/Makefile | 43 + samples/2_Cookbook/13_occupancy/occupancy.cpp | 177 + samples/2_Cookbook/14_gpu_arch/CMakeLists.txt | 44 + samples/2_Cookbook/14_gpu_arch/Makefile | 43 + samples/2_Cookbook/14_gpu_arch/gpuarch.cpp | 95 + .../2_Cookbook/15_static_library/README.md | 179 + .../device_functions/CMakeLists.txt | 58 + .../device_functions/Makefile | 35 + .../device_functions/hipDevice.cpp | 28 + .../device_functions/hipMain2.cpp | 69 + .../host_functions/CMakeLists.txt | 55 + .../15_static_library/host_functions/Makefile | 42 + .../host_functions/hipMain1.cpp | 28 + .../host_functions/hipOptLibrary.cpp | 62 + .../16_assembly_to_executable/Makefile | 82 + .../16_assembly_to_executable/README.md | 53 + .../hip_obj_gen.mcin | 20 + .../16_assembly_to_executable/square.cpp | 96 + .../17_llvm_ir_to_executable/Makefile | 102 + .../17_llvm_ir_to_executable/README.md | 68 + .../17_llvm_ir_to_executable/hip_obj_gen.mcin | 20 + .../17_llvm_ir_to_executable/square.cpp | 96 + .../18_cmake_hip_device/CMakeLists.txt | 32 + .../2_Cookbook/18_cmake_hip_device/README.md | 16 + .../2_Cookbook/18_cmake_hip_device/square.cpp | 98 + .../2_Cookbook/19_cmake_lang/CMakeLists.txt | 14 + .../19_cmake_lang/MatrixTranspose.cpp | 122 + samples/2_Cookbook/19_cmake_lang/README.md | 20 + .../2_Cookbook/19_cmake_lang/TestFortran.F90 | 3 + samples/2_Cookbook/1_hipEvent/CMakeLists.txt | 44 + samples/2_Cookbook/1_hipEvent/Makefile | 57 + samples/2_Cookbook/1_hipEvent/Readme.md | 80 + samples/2_Cookbook/1_hipEvent/hipEvent.cpp | 158 + .../2_Cookbook/20_hip_vulkan/CMakeLists.txt | 97 + .../20_hip_vulkan/SineWaveSimulation.cpp | 147 + .../20_hip_vulkan/SineWaveSimulation.h | 101 + .../20_hip_vulkan/SineWaveSimulation.hip | 147 + .../20_hip_vulkan/VulkanBaseApp.cpp | 1724 ++ .../2_Cookbook/20_hip_vulkan/VulkanBaseApp.h | 146 + samples/2_Cookbook/20_hip_vulkan/buildcmd.txt | 16 + samples/2_Cookbook/20_hip_vulkan/linmath.h | 502 + samples/2_Cookbook/20_hip_vulkan/main.cpp | 466 + .../2_Cookbook/20_hip_vulkan/sinewave.frag | 38 + .../2_Cookbook/20_hip_vulkan/sinewave.vert | 43 + .../21_cmake_hip_cxx_clang/CMakeLists.txt | 34 + .../21_cmake_hip_cxx_clang/README.md | 19 + .../21_cmake_hip_cxx_clang/square.cpp | 98 + .../22_cmake_hip_lang/CMakeLists.txt | 27 + .../2_Cookbook/22_cmake_hip_lang/README.md | 16 + .../2_Cookbook/22_cmake_hip_lang/square.hip | 98 + .../2_Cookbook/23_cmake_hiprtc/CMakeLists.txt | 36 + samples/2_Cookbook/23_cmake_hiprtc/README.md | 9 + samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp | 156 + .../2_Cookbook/3_shared_memory/CMakeLists.txt | 44 + samples/2_Cookbook/3_shared_memory/Makefile | 59 + samples/2_Cookbook/3_shared_memory/Readme.md | 42 + .../3_shared_memory/sharedMemory.cpp | 128 + samples/2_Cookbook/4_shfl/CMakeLists.txt | 44 + samples/2_Cookbook/4_shfl/Makefile | 62 + samples/2_Cookbook/4_shfl/Readme.md | 53 + samples/2_Cookbook/4_shfl/shfl.cpp | 124 + samples/2_Cookbook/5_2dshfl/2dshfl.cpp | 122 + samples/2_Cookbook/5_2dshfl/CMakeLists.txt | 43 + samples/2_Cookbook/5_2dshfl/Makefile | 63 + samples/2_Cookbook/5_2dshfl/Readme.md | 55 + .../6_dynamic_shared/CMakeLists.txt | 43 + samples/2_Cookbook/6_dynamic_shared/Makefile | 59 + samples/2_Cookbook/6_dynamic_shared/Readme.md | 52 + .../6_dynamic_shared/dynamic_shared.cpp | 127 + samples/2_Cookbook/7_streams/CMakeLists.txt | 43 + samples/2_Cookbook/7_streams/Makefile | 59 + samples/2_Cookbook/7_streams/Readme.md | 63 + samples/2_Cookbook/7_streams/stream.cpp | 138 + samples/2_Cookbook/8_peer2peer/CMakeLists.txt | 43 + samples/2_Cookbook/8_peer2peer/Makefile | 59 + samples/2_Cookbook/8_peer2peer/peer2peer.cpp | 226 + samples/2_Cookbook/9_unroll/CMakeLists.txt | 43 + samples/2_Cookbook/9_unroll/Makefile | 62 + samples/2_Cookbook/9_unroll/Readme.md | 48 + samples/2_Cookbook/9_unroll/unroll.cpp | 122 + samples/README.md | 33 + 592 files changed, 109310 insertions(+) create mode 100644 LICENSE.txt create mode 100644 catch/ABM/AddKernels/CMakeLists.txt create mode 100644 catch/ABM/AddKernels/add.cc create mode 100644 catch/ABM/CMakeLists.txt create mode 100644 catch/CMakeLists.txt create mode 100644 catch/README.md create mode 100644 catch/TypeQualifiers/CMakeLists.txt create mode 100644 catch/TypeQualifiers/hipManagedKeyword.cc create mode 100644 catch/catchProp_in_rc.in create mode 100644 catch/external/Catch2/LICENSE.txt create mode 100644 catch/external/Catch2/catch.hpp create mode 100644 catch/external/Catch2/cmake/Catch2/Catch.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/Catch2Config.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/Catch2ConfigVersion.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/Catch2Targets.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/CatchAddTests.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/ParseAndAddCatchTests.cmake create mode 100644 catch/external/Catch2/cmake/Catch2/catch_include.cmake.in create mode 100644 catch/external/picojson/LICENSE create mode 100644 catch/external/picojson/picojson.h create mode 100644 catch/hipTestMain/CMakeLists.txt create mode 100644 catch/hipTestMain/config/config_amd_linux_MI2xx.json create mode 100644 catch/hipTestMain/config/config_amd_linux_common.json create mode 100644 catch/hipTestMain/config/config_amd_windows_MI2xx.json create mode 100644 catch/hipTestMain/config/config_amd_windows_common.json create mode 100644 catch/hipTestMain/hip_test_context.cc create mode 100644 catch/hipTestMain/main.cc create mode 100644 catch/hipTestMain/standalone_main.cc create mode 100644 catch/include/hip_test_checkers.hh create mode 100644 catch/include/hip_test_common.hh create mode 100644 catch/include/hip_test_context.hh create mode 100644 catch/include/hip_test_filesystem.hh create mode 100644 catch/include/hip_test_helper.hh create mode 100644 catch/include/hip_test_kernels.hh create mode 100644 catch/include/hip_test_process.hh create mode 100644 catch/include/hip_test_rtc.hh create mode 100644 catch/include/hip_test_smi.hh create mode 100644 catch/include/hip_texture_helper.hh create mode 100644 catch/include/kernel_mapping.hh create mode 100644 catch/include/kernels.hh create mode 100644 catch/include/resource_guards.hh create mode 100644 catch/include/threaded_zig_zag_test.hh create mode 100644 catch/include/utils.hh create mode 100644 catch/kernels/CMakeLists.txt create mode 100644 catch/kernels/Set.cpp create mode 100644 catch/kernels/vectorADD.inl create mode 100644 catch/multiproc/CMakeLists.txt create mode 100644 catch/multiproc/childMalloc.cc create mode 100644 catch/multiproc/deviceAllocationMproc.cc create mode 100644 catch/multiproc/dummy_kernel.cpp create mode 100644 catch/multiproc/hipDeviceComputeCapabilityMproc.cc create mode 100644 catch/multiproc/hipDeviceGetPCIBusIdMproc.cc create mode 100644 catch/multiproc/hipDeviceTotalMemMproc.cc create mode 100644 catch/multiproc/hipGetDeviceAttributeMproc.cc create mode 100644 catch/multiproc/hipGetDeviceCountMproc.cc create mode 100644 catch/multiproc/hipGetDevicePropertiesMproc.cc create mode 100644 catch/multiproc/hipIpcEventHandle.cc create mode 100644 catch/multiproc/hipIpcMemAccessTest.cc create mode 100644 catch/multiproc/hipMallocConcurrencyMproc.cc create mode 100644 catch/multiproc/hipMemCoherencyTstMProc.cc create mode 100644 catch/multiproc/hipNoGpuTsts.cc create mode 100644 catch/multiproc/hipSetGetDeviceMproc.cc create mode 100644 catch/packaging/hip-tests.txt create mode 100644 catch/stress/CMakeLists.txt create mode 100644 catch/stress/deviceallocation/CMakeLists.txt create mode 100644 catch/stress/deviceallocation/Stress_deviceAllocationStress.cc create mode 100644 catch/stress/memory/CMakeLists.txt create mode 100644 catch/stress/memory/hipHostMalloc.cc create mode 100644 catch/stress/memory/hipMalloc.cc create mode 100644 catch/stress/memory/hipMallocManagedStress.cc create mode 100644 catch/stress/memory/hipMemPrftchAsyncStressTst.cc create mode 100644 catch/stress/memory/hipMemcpyMThreadMSize.cc create mode 100644 catch/stress/memory/memcpy.cc create mode 100644 catch/stress/printf/CMakeLists.txt create mode 100644 catch/stress/printf/Stress_printf_ComplexKernels.cc create mode 100644 catch/stress/printf/Stress_printf_SimpleKernels.cc create mode 100644 catch/stress/printf/printf_common.h create mode 100644 catch/stress/stream/CMakeLists.txt create mode 100644 catch/stress/stream/Stress_hipStreamCreate.cc create mode 100644 catch/stress/stream/streamEnqueue.cc create mode 100644 catch/unit/CMakeLists.txt create mode 100644 catch/unit/clock/CMakeLists.txt create mode 100644 catch/unit/clock/hipClockCheck.cc create mode 100644 catch/unit/cooperativeGrps/coalesced_groups_shfl_down.cc create mode 100644 catch/unit/cooperativeGrps/coalesced_groups_shfl_up.cc create mode 100644 catch/unit/cooperativeGrps/simple_coalesced_groups.cc create mode 100644 catch/unit/device/CMakeLists.txt create mode 100644 catch/unit/device/getDeviceCount_exe.cc create mode 100644 catch/unit/device/hipChooseDevice.cc create mode 100644 catch/unit/device/hipDeviceComputeCapability.cc create mode 100644 catch/unit/device/hipDeviceGetByPCIBusId.cc create mode 100644 catch/unit/device/hipDeviceGetLimit.cc create mode 100644 catch/unit/device/hipDeviceGetName.cc create mode 100644 catch/unit/device/hipDeviceGetP2PAttribute.cc create mode 100644 catch/unit/device/hipDeviceGetP2PAttribute_exe.cc create mode 100644 catch/unit/device/hipDeviceGetPCIBusId.cc create mode 100644 catch/unit/device/hipDeviceGetUuid.cc create mode 100644 catch/unit/device/hipDeviceSetGetCacheConfig.cc create mode 100644 catch/unit/device/hipDeviceSetLimit.cc create mode 100644 catch/unit/device/hipDeviceSynchronize.cc create mode 100644 catch/unit/device/hipDeviceTotalMem.cc create mode 100644 catch/unit/device/hipExtGetLinkTypeAndHopCount.cc create mode 100644 catch/unit/device/hipGetDeviceAttribute.cc create mode 100644 catch/unit/device/hipGetDeviceCount.cc create mode 100644 catch/unit/device/hipGetDeviceProperties.cc create mode 100644 catch/unit/device/hipGetSetDeviceFlags.cc create mode 100644 catch/unit/device/hipRuntimeGetVersion.cc create mode 100644 catch/unit/device/hipSetGetDevice.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc create mode 100644 catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc create mode 100644 catch/unit/deviceLib/BuiltIns_fadd.cc create mode 100644 catch/unit/deviceLib/BuiltIns_fmax.cc create mode 100644 catch/unit/deviceLib/BuiltIns_fmin.cc create mode 100644 catch/unit/deviceLib/CMakeLists.txt create mode 100644 catch/unit/deviceLib/anyAll.cc create mode 100644 catch/unit/deviceLib/ballot.cc create mode 100644 catch/unit/deviceLib/bitExtract.cc create mode 100644 catch/unit/deviceLib/bitInsert.cc create mode 100644 catch/unit/deviceLib/brev.cc create mode 100644 catch/unit/deviceLib/clz.cc create mode 100644 catch/unit/deviceLib/defs.h create mode 100644 catch/unit/deviceLib/deviceAllocCommon.h create mode 100644 catch/unit/deviceLib/deviceAllocation.cc create mode 100644 catch/unit/deviceLib/ffs.cc create mode 100644 catch/unit/deviceLib/floatMath.cc create mode 100644 catch/unit/deviceLib/floatTM.cc create mode 100644 catch/unit/deviceLib/funnelshift.cc create mode 100644 catch/unit/deviceLib/hipTestDeviceSymbol.cc create mode 100644 catch/unit/deviceLib/kerDevAllocMultCO.cc create mode 100644 catch/unit/deviceLib/kerDevAllocSingleKer.cc create mode 100644 catch/unit/deviceLib/kerDevFreeMultCO.cc create mode 100644 catch/unit/deviceLib/kerDevWriteMultCO.cc create mode 100644 catch/unit/deviceLib/ldg.cc create mode 100644 catch/unit/deviceLib/mbcnt.cc create mode 100644 catch/unit/deviceLib/popc.cc create mode 100644 catch/unit/deviceLib/syncthreadsand.cc create mode 100644 catch/unit/deviceLib/syncthreadscount.cc create mode 100644 catch/unit/deviceLib/syncthreadsor.cc create mode 100644 catch/unit/deviceLib/threadfence_system.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc create mode 100644 catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc create mode 100644 catch/unit/deviceLib/vectorTypesDevice.cc create mode 100644 catch/unit/event/CMakeLists.txt create mode 100644 catch/unit/event/Unit_hipEvent.cc create mode 100644 catch/unit/event/Unit_hipEventElapsedTime.cc create mode 100644 catch/unit/event/Unit_hipEventIpc.cc create mode 100644 catch/unit/event/Unit_hipEventQuery.cc create mode 100644 catch/unit/event/Unit_hipEventRecord.cc create mode 100644 catch/unit/event/Unit_hipEvent_Negative.cc create mode 100644 catch/unit/event/hipEventCreateWithFlags.cc create mode 100644 catch/unit/event/hipEventDestroy.cc create mode 100644 catch/unit/event/hipEventSynchronize.cc create mode 100644 catch/unit/graph/CMakeLists.txt create mode 100644 catch/unit/graph/hipGraph.cc create mode 100644 catch/unit/graph/hipGraphAddChildGraphNode.cc create mode 100644 catch/unit/graph/hipGraphAddDependencies.cc create mode 100644 catch/unit/graph/hipGraphAddEmptyNode.cc create mode 100644 catch/unit/graph/hipGraphAddEventRecordNode.cc create mode 100644 catch/unit/graph/hipGraphAddHostNode.cc create mode 100644 catch/unit/graph/hipGraphAddKernelNode.cc create mode 100644 catch/unit/graph/hipGraphAddMemcpyNode.cc create mode 100644 catch/unit/graph/hipGraphAddMemcpyNode1D.cc create mode 100644 catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc create mode 100644 catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc create mode 100644 catch/unit/graph/hipGraphAddMemsetNode.cc create mode 100644 catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc create mode 100644 catch/unit/graph/hipGraphClone.cc create mode 100644 catch/unit/graph/hipGraphDestroyNode.cc create mode 100644 catch/unit/graph/hipGraphEventRecordNodeGetEvent.cc create mode 100644 catch/unit/graph/hipGraphEventRecordNodeSetEvent.cc create mode 100644 catch/unit/graph/hipGraphEventWaitNodeGetEvent.cc create mode 100644 catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc create mode 100644 catch/unit/graph/hipGraphExecChildGraphNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc create mode 100644 catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc create mode 100644 catch/unit/graph/hipGraphExecHostNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecKernelNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecMemcpyNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc create mode 100644 catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc create mode 100644 catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc create mode 100644 catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphExecUpdate.cc create mode 100644 catch/unit/graph/hipGraphGetEdges.cc create mode 100644 catch/unit/graph/hipGraphGetNodes.cc create mode 100644 catch/unit/graph/hipGraphGetRootNodes.cc create mode 100644 catch/unit/graph/hipGraphHostNodeGetParams.cc create mode 100644 catch/unit/graph/hipGraphHostNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphInstantiate.cc create mode 100644 catch/unit/graph/hipGraphInstantiateWithFlags.cc create mode 100644 catch/unit/graph/hipGraphKernelNodeGetParams.cc create mode 100644 catch/unit/graph/hipGraphKernelNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphLaunch.cc create mode 100644 catch/unit/graph/hipGraphMemcpyNodeGetParams.cc create mode 100644 catch/unit/graph/hipGraphMemcpyNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc create mode 100644 catch/unit/graph/hipGraphMemcpyNodeSetParamsFromSymbol.cc create mode 100644 catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc create mode 100644 catch/unit/graph/hipGraphMemsetNodeGetParams.cc create mode 100644 catch/unit/graph/hipGraphMemsetNodeSetParams.cc create mode 100644 catch/unit/graph/hipGraphNodeFindInClone.cc create mode 100644 catch/unit/graph/hipGraphNodeGetDependencies.cc create mode 100644 catch/unit/graph/hipGraphNodeGetDependentNodes.cc create mode 100644 catch/unit/graph/hipGraphNodeGetType.cc create mode 100644 catch/unit/graph/hipGraphRemoveDependencies.cc create mode 100644 catch/unit/graph/hipSimpleGraphWithKernel.cc create mode 100644 catch/unit/graph/hipStreamBeginCapture.cc create mode 100644 catch/unit/graph/hipStreamEndCapture.cc create mode 100644 catch/unit/graph/hipStreamGetCaptureInfo.cc create mode 100644 catch/unit/graph/hipStreamIsCapturing.cc create mode 100644 catch/unit/kernel/CMakeLists.txt create mode 100644 catch/unit/kernel/hipLaunchBounds.cc create mode 100644 catch/unit/kernel/hipMemFaultStackAllocation.cc create mode 100644 catch/unit/memory/CMakeLists.txt create mode 100644 catch/unit/memory/DriverContext.cc create mode 100644 catch/unit/memory/DriverContext.hh create mode 100644 catch/unit/memory/MemUtils.hh create mode 100644 catch/unit/memory/hipArray.cc create mode 100644 catch/unit/memory/hipArray3DCreate.cc create mode 100644 catch/unit/memory/hipArrayCommon.hh create mode 100644 catch/unit/memory/hipArrayCreate.cc create mode 100644 catch/unit/memory/hipDrvMemcpy3D.cc create mode 100644 catch/unit/memory/hipDrvMemcpy3DAsync.cc create mode 100644 catch/unit/memory/hipDrvPtrGetAttributes.cc create mode 100644 catch/unit/memory/hipFree.cc create mode 100644 catch/unit/memory/hipHmmOvrSubscriptionTst.cc create mode 100644 catch/unit/memory/hipHostGetDevicePointer.cc create mode 100644 catch/unit/memory/hipHostGetFlags.cc create mode 100644 catch/unit/memory/hipHostMalloc.cc create mode 100644 catch/unit/memory/hipHostMallocTests.cc create mode 100644 catch/unit/memory/hipHostRegister.cc create mode 100644 catch/unit/memory/hipHostUnregister.cc create mode 100644 catch/unit/memory/hipMalloc3D.cc create mode 100644 catch/unit/memory/hipMalloc3DArray.cc create mode 100644 catch/unit/memory/hipMallocArray.cc create mode 100644 catch/unit/memory/hipMallocConcurrency.cc create mode 100644 catch/unit/memory/hipMallocManaged.cc create mode 100644 catch/unit/memory/hipMallocManagedCommon.hh create mode 100644 catch/unit/memory/hipMallocManagedFlagsTst.cc create mode 100644 catch/unit/memory/hipMallocManaged_MultiScenario.cc create mode 100644 catch/unit/memory/hipMallocMngdMultiThread.cc create mode 100644 catch/unit/memory/hipMallocPitch.cc create mode 100644 catch/unit/memory/hipMemAdvise.cc create mode 100644 catch/unit/memory/hipMemAdviseMmap.cc create mode 100644 catch/unit/memory/hipMemCoherencyTst.cc create mode 100644 catch/unit/memory/hipMemGetInfo.cc create mode 100644 catch/unit/memory/hipMemPoolApi.cc create mode 100644 catch/unit/memory/hipMemPrefetchAsync.cc create mode 100644 catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc create mode 100644 catch/unit/memory/hipMemPtrGetInfo.cc create mode 100644 catch/unit/memory/hipMemRangeGetAttribute.cc create mode 100644 catch/unit/memory/hipMemVmm.cc create mode 100644 catch/unit/memory/hipMemcpy.cc create mode 100644 catch/unit/memory/hipMemcpy2D.cc create mode 100644 catch/unit/memory/hipMemcpy2DAsync.cc create mode 100644 catch/unit/memory/hipMemcpy2DFromArray.cc create mode 100644 catch/unit/memory/hipMemcpy2DFromArrayAsync.cc create mode 100644 catch/unit/memory/hipMemcpy2DToArray.cc create mode 100644 catch/unit/memory/hipMemcpy2DToArrayAsync.cc create mode 100644 catch/unit/memory/hipMemcpy3D.cc create mode 100644 catch/unit/memory/hipMemcpy3DAsync.cc create mode 100644 catch/unit/memory/hipMemcpyAllApiNegative.cc create mode 100644 catch/unit/memory/hipMemcpyAsync.cc create mode 100644 catch/unit/memory/hipMemcpyAtoH.cc create mode 100644 catch/unit/memory/hipMemcpyDtoD.cc create mode 100644 catch/unit/memory/hipMemcpyDtoDAsync.cc create mode 100644 catch/unit/memory/hipMemcpyFromSymbol.cc create mode 100644 catch/unit/memory/hipMemcpyHtoA.cc create mode 100644 catch/unit/memory/hipMemcpyParam2D.cc create mode 100644 catch/unit/memory/hipMemcpyParam2DAsync.cc create mode 100644 catch/unit/memory/hipMemcpyPeer.cc create mode 100644 catch/unit/memory/hipMemcpyPeerAsync.cc create mode 100644 catch/unit/memory/hipMemcpySync.cc create mode 100644 catch/unit/memory/hipMemcpyWithStream.cc create mode 100644 catch/unit/memory/hipMemcpyWithStreamMultiThread.cc create mode 100644 catch/unit/memory/hipMemcpy_MultiThread.cc create mode 100644 catch/unit/memory/hipMemoryAllocateCoherent.cc create mode 100644 catch/unit/memory/hipMemset.cc create mode 100644 catch/unit/memory/hipMemset2D.cc create mode 100644 catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc create mode 100644 catch/unit/memory/hipMemset3D.cc create mode 100644 catch/unit/memory/hipMemset3DFunctional.cc create mode 100644 catch/unit/memory/hipMemset3DRegressMultiThread.cc create mode 100644 catch/unit/memory/hipMemsetAsync.cc create mode 100644 catch/unit/memory/hipMemsetAsyncAndKernel.cc create mode 100644 catch/unit/memory/hipMemsetAsyncMultiThread.cc create mode 100644 catch/unit/memory/hipMemsetFunctional.cc create mode 100644 catch/unit/memory/hipMemsetNegative.cc create mode 100644 catch/unit/memory/hipMemsetSync.cc create mode 100644 catch/unit/memory/hipPointerGetAttribute.cc create mode 100644 catch/unit/memory/hipPointerGetAttributes.cc create mode 100644 catch/unit/memory/hipPtrGetAttribute.cc create mode 100644 catch/unit/memory/malloc.cc create mode 100644 catch/unit/memory/memset.cc create mode 100644 catch/unit/multiThread/CMakeLists.txt create mode 100644 catch/unit/multiThread/hipMultiThreadDevice.cc create mode 100644 catch/unit/multiThread/hipMultiThreadStreams1.cc create mode 100644 catch/unit/multiThread/hipMultiThreadStreams2.cc create mode 100644 catch/unit/occupancy/CMakeLists.txt create mode 100644 catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc create mode 100644 catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc create mode 100644 catch/unit/printf/CMakeLists.txt create mode 100644 catch/unit/printf/printfFlags.cc create mode 100644 catch/unit/printf/printfFlags_exe.cc create mode 100644 catch/unit/printf/printfSpecifiers.cc create mode 100644 catch/unit/printf/printfSpecifiers_exe.cc create mode 100644 catch/unit/rtc/CMakeLists.txt create mode 100644 catch/unit/rtc/customOptions.cc create mode 100644 catch/unit/rtc/headers/test_header1.h create mode 100644 catch/unit/rtc/headers/test_header2.h create mode 100644 catch/unit/rtc/hipRtcFunctional.cc create mode 100644 catch/unit/rtc/includepath.cc create mode 100644 catch/unit/rtc/saxpy.cc create mode 100644 catch/unit/rtc/saxpy.h create mode 100644 catch/unit/rtc/warpsize.cc create mode 100644 catch/unit/stream/CMakeLists.txt create mode 100644 catch/unit/stream/hipAPIStreamDisable.cc create mode 100644 catch/unit/stream/hipDeviceGetStreamPriorityRange.cc create mode 100644 catch/unit/stream/hipMultiStream.cc create mode 100644 catch/unit/stream/hipStreamACb_MultiThread.cc create mode 100644 catch/unit/stream/hipStreamAddCallback.cc create mode 100644 catch/unit/stream/hipStreamAttachMemAsync.cc create mode 100644 catch/unit/stream/hipStreamCreate.cc create mode 100644 catch/unit/stream/hipStreamCreateWithFlags.cc create mode 100644 catch/unit/stream/hipStreamCreateWithPriority.cc create mode 100644 catch/unit/stream/hipStreamDestroy.cc create mode 100644 catch/unit/stream/hipStreamGetCUMask.cc create mode 100644 catch/unit/stream/hipStreamGetFlags.cc create mode 100644 catch/unit/stream/hipStreamGetPriority.cc create mode 100644 catch/unit/stream/hipStreamQuery.cc create mode 100644 catch/unit/stream/hipStreamSynchronize.cc create mode 100644 catch/unit/stream/hipStreamValue.cc create mode 100644 catch/unit/stream/hipStreamWaitEvent.cc create mode 100644 catch/unit/stream/hipStreamWithCUMask.cc create mode 100644 catch/unit/stream/streamCommon.cc create mode 100644 catch/unit/stream/streamCommon.hh create mode 100644 catch/unit/streamperthread/CMakeLists.txt create mode 100644 catch/unit/streamperthread/hipStreamPerThrdTsts.cc create mode 100644 catch/unit/streamperthread/hipStreamPerThread_Basic.cc create mode 100644 catch/unit/streamperthread/hipStreamPerThread_DeviceReset.cc create mode 100644 catch/unit/streamperthread/hipStreamPerThread_Event.cc create mode 100644 catch/unit/streamperthread/hipStreamPerThread_MultiThread.cc create mode 100644 catch/unit/texture/CMakeLists.txt create mode 100644 catch/unit/texture/hipBindTex2DPitch.cc create mode 100644 catch/unit/texture/hipBindTexRef1DFetch.cc create mode 100644 catch/unit/texture/hipCreateTextureObject_ArgValidation.cc create mode 100644 catch/unit/texture/hipCreateTextureObject_Array.cc create mode 100644 catch/unit/texture/hipCreateTextureObject_Linear.cc create mode 100644 catch/unit/texture/hipCreateTextureObject_Pitch2D.cc create mode 100644 catch/unit/texture/hipGetChanDesc.cc create mode 100644 catch/unit/texture/hipNormalizedFloatValueTex.cc create mode 100644 catch/unit/texture/hipSimpleTexture2DLayered.cc create mode 100644 catch/unit/texture/hipSimpleTexture3D.cc create mode 100644 catch/unit/texture/hipTex1DFetchCheckModes.cc create mode 100644 catch/unit/texture/hipTexObjPitch.cc create mode 100644 catch/unit/texture/hipTextureMipmapObj2D.cc create mode 100644 catch/unit/texture/hipTextureObj1DCheckModes.cc create mode 100644 catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc create mode 100644 catch/unit/texture/hipTextureObj1DFetch.cc create mode 100644 catch/unit/texture/hipTextureObj2D.cc create mode 100644 catch/unit/texture/hipTextureObj2DCheckModes.cc create mode 100644 catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc create mode 100644 catch/unit/texture/hipTextureObj3DCheckModes.cc create mode 100644 catch/unit/texture/hipTextureObjFetchVector.cc create mode 100644 catch/unit/texture/hipTextureRef2D.cc create mode 100644 samples/0_Intro/bit_extract/CMakeLists.txt create mode 100644 samples/0_Intro/bit_extract/Makefile create mode 100644 samples/0_Intro/bit_extract/README.md create mode 100644 samples/0_Intro/bit_extract/bit_extract.cpp create mode 100644 samples/0_Intro/module_api/CMakeLists.txt create mode 100644 samples/0_Intro/module_api/Makefile create mode 100644 samples/0_Intro/module_api/defaultDriver.cpp create mode 100644 samples/0_Intro/module_api/launchKernelHcc.cpp create mode 100644 samples/0_Intro/module_api/runKernel.cpp create mode 100644 samples/0_Intro/module_api/vcpy_kernel.cpp create mode 100644 samples/0_Intro/module_api_global/CMakeLists.txt create mode 100644 samples/0_Intro/module_api_global/Makefile create mode 100644 samples/0_Intro/module_api_global/runKernel.cpp create mode 100644 samples/0_Intro/module_api_global/vcpy_kernel.cpp create mode 100644 samples/0_Intro/square/CMakeLists.txt create mode 100644 samples/0_Intro/square/Makefile create mode 100644 samples/0_Intro/square/README.md create mode 100644 samples/0_Intro/square/square.cu create mode 100644 samples/0_Intro/square/square.hipref.cpp create mode 100644 samples/1_Utils/hipBusBandwidth/CMakeLists.txt create mode 100644 samples/1_Utils/hipBusBandwidth/LICENSE.txt create mode 100644 samples/1_Utils/hipBusBandwidth/Makefile create mode 100644 samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp create mode 100644 samples/1_Utils/hipBusBandwidth/ResultDatabase.h create mode 100644 samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp create mode 100644 samples/1_Utils/hipCommander/CMakeLists.txt create mode 100644 samples/1_Utils/hipCommander/LICENSE.txt create mode 100644 samples/1_Utils/hipCommander/Makefile create mode 100644 samples/1_Utils/hipCommander/ResultDatabase.cpp create mode 100644 samples/1_Utils/hipCommander/ResultDatabase.h create mode 100644 samples/1_Utils/hipCommander/c.cmd create mode 100644 samples/1_Utils/hipCommander/classic.cmd create mode 100644 samples/1_Utils/hipCommander/hipCommander.cpp create mode 100644 samples/1_Utils/hipCommander/l2.hcm create mode 100644 samples/1_Utils/hipCommander/loop.hcm create mode 100644 samples/1_Utils/hipCommander/loop2.hcm create mode 100644 samples/1_Utils/hipCommander/nullkernel.hip.cpp create mode 100755 samples/1_Utils/hipCommander/nullkernel.hsaco create mode 100644 samples/1_Utils/hipCommander/perf/latency2.hcm create mode 100644 samples/1_Utils/hipCommander/perf/latency_hostsync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/latency_nosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/latency_nullstream.hcm create mode 100644 samples/1_Utils/hipCommander/perf/modulelaunch_latency.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm create mode 100644 samples/1_Utils/hipCommander/setstream.hcm create mode 100644 samples/1_Utils/hipCommander/testcase.cpp create mode 100644 samples/1_Utils/hipDispatchLatency/CMakeLists.txt create mode 100644 samples/1_Utils/hipDispatchLatency/Makefile create mode 100644 samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp create mode 100644 samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp create mode 100644 samples/1_Utils/hipDispatchLatency/test_kernel.cpp create mode 100644 samples/1_Utils/hipInfo/CMakeLists.txt create mode 100644 samples/1_Utils/hipInfo/Makefile create mode 100644 samples/1_Utils/hipInfo/README.md create mode 100644 samples/1_Utils/hipInfo/hipInfo.cpp create mode 100644 samples/2_Cookbook/0_MatrixTranspose/CMakeLists.txt create mode 100644 samples/2_Cookbook/0_MatrixTranspose/Makefile create mode 100644 samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp create mode 100644 samples/2_Cookbook/0_MatrixTranspose/Readme.md create mode 100644 samples/2_Cookbook/10_inline_asm/CMakeLists.txt create mode 100644 samples/2_Cookbook/10_inline_asm/Makefile create mode 100644 samples/2_Cookbook/10_inline_asm/Readme.md create mode 100644 samples/2_Cookbook/10_inline_asm/inline_asm.cpp create mode 100644 samples/2_Cookbook/11_texture_driver/CMakeLists.txt create mode 100644 samples/2_Cookbook/11_texture_driver/Makefile create mode 100644 samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp create mode 100644 samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp create mode 100644 samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt create mode 100644 samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp create mode 100644 samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md create mode 100644 samples/2_Cookbook/13_occupancy/CMakeLists.txt create mode 100644 samples/2_Cookbook/13_occupancy/Makefile create mode 100644 samples/2_Cookbook/13_occupancy/occupancy.cpp create mode 100644 samples/2_Cookbook/14_gpu_arch/CMakeLists.txt create mode 100644 samples/2_Cookbook/14_gpu_arch/Makefile create mode 100644 samples/2_Cookbook/14_gpu_arch/gpuarch.cpp create mode 100644 samples/2_Cookbook/15_static_library/README.md create mode 100644 samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt create mode 100644 samples/2_Cookbook/15_static_library/device_functions/Makefile create mode 100644 samples/2_Cookbook/15_static_library/device_functions/hipDevice.cpp create mode 100644 samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp create mode 100644 samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt create mode 100644 samples/2_Cookbook/15_static_library/host_functions/Makefile create mode 100644 samples/2_Cookbook/15_static_library/host_functions/hipMain1.cpp create mode 100644 samples/2_Cookbook/15_static_library/host_functions/hipOptLibrary.cpp create mode 100644 samples/2_Cookbook/16_assembly_to_executable/Makefile create mode 100644 samples/2_Cookbook/16_assembly_to_executable/README.md create mode 100644 samples/2_Cookbook/16_assembly_to_executable/hip_obj_gen.mcin create mode 100644 samples/2_Cookbook/16_assembly_to_executable/square.cpp create mode 100644 samples/2_Cookbook/17_llvm_ir_to_executable/Makefile create mode 100644 samples/2_Cookbook/17_llvm_ir_to_executable/README.md create mode 100644 samples/2_Cookbook/17_llvm_ir_to_executable/hip_obj_gen.mcin create mode 100644 samples/2_Cookbook/17_llvm_ir_to_executable/square.cpp create mode 100644 samples/2_Cookbook/18_cmake_hip_device/CMakeLists.txt create mode 100644 samples/2_Cookbook/18_cmake_hip_device/README.md create mode 100644 samples/2_Cookbook/18_cmake_hip_device/square.cpp create mode 100644 samples/2_Cookbook/19_cmake_lang/CMakeLists.txt create mode 100644 samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp create mode 100644 samples/2_Cookbook/19_cmake_lang/README.md create mode 100644 samples/2_Cookbook/19_cmake_lang/TestFortran.F90 create mode 100644 samples/2_Cookbook/1_hipEvent/CMakeLists.txt create mode 100644 samples/2_Cookbook/1_hipEvent/Makefile create mode 100644 samples/2_Cookbook/1_hipEvent/Readme.md create mode 100644 samples/2_Cookbook/1_hipEvent/hipEvent.cpp create mode 100644 samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt create mode 100644 samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp create mode 100644 samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.h create mode 100644 samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.hip create mode 100644 samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.cpp create mode 100644 samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.h create mode 100644 samples/2_Cookbook/20_hip_vulkan/buildcmd.txt create mode 100644 samples/2_Cookbook/20_hip_vulkan/linmath.h create mode 100644 samples/2_Cookbook/20_hip_vulkan/main.cpp create mode 100644 samples/2_Cookbook/20_hip_vulkan/sinewave.frag create mode 100644 samples/2_Cookbook/20_hip_vulkan/sinewave.vert create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/README.md create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/square.hip create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/CMakeLists.txt create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/README.md create mode 100644 samples/2_Cookbook/23_cmake_hiprtc/saxpy.cpp create mode 100644 samples/2_Cookbook/3_shared_memory/CMakeLists.txt create mode 100644 samples/2_Cookbook/3_shared_memory/Makefile create mode 100644 samples/2_Cookbook/3_shared_memory/Readme.md create mode 100644 samples/2_Cookbook/3_shared_memory/sharedMemory.cpp create mode 100644 samples/2_Cookbook/4_shfl/CMakeLists.txt create mode 100644 samples/2_Cookbook/4_shfl/Makefile create mode 100644 samples/2_Cookbook/4_shfl/Readme.md create mode 100644 samples/2_Cookbook/4_shfl/shfl.cpp create mode 100644 samples/2_Cookbook/5_2dshfl/2dshfl.cpp create mode 100644 samples/2_Cookbook/5_2dshfl/CMakeLists.txt create mode 100644 samples/2_Cookbook/5_2dshfl/Makefile create mode 100644 samples/2_Cookbook/5_2dshfl/Readme.md create mode 100644 samples/2_Cookbook/6_dynamic_shared/CMakeLists.txt create mode 100644 samples/2_Cookbook/6_dynamic_shared/Makefile create mode 100644 samples/2_Cookbook/6_dynamic_shared/Readme.md create mode 100644 samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp create mode 100644 samples/2_Cookbook/7_streams/CMakeLists.txt create mode 100644 samples/2_Cookbook/7_streams/Makefile create mode 100644 samples/2_Cookbook/7_streams/Readme.md create mode 100644 samples/2_Cookbook/7_streams/stream.cpp create mode 100644 samples/2_Cookbook/8_peer2peer/CMakeLists.txt create mode 100644 samples/2_Cookbook/8_peer2peer/Makefile create mode 100644 samples/2_Cookbook/8_peer2peer/peer2peer.cpp create mode 100644 samples/2_Cookbook/9_unroll/CMakeLists.txt create mode 100644 samples/2_Cookbook/9_unroll/Makefile create mode 100644 samples/2_Cookbook/9_unroll/Readme.md create mode 100644 samples/2_Cookbook/9_unroll/unroll.cpp create mode 100644 samples/README.md diff --git a/LICENSE.txt b/LICENSE.txt new file mode 100644 index 000000000..4cbb63923 --- /dev/null +++ b/LICENSE.txt @@ -0,0 +1,20 @@ +Copyright (c) 2008 - 2022 Advanced Micro Devices, Inc. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. + diff --git a/catch/ABM/AddKernels/CMakeLists.txt b/catch/ABM/AddKernels/CMakeLists.txt new file mode 100644 index 000000000..776f0fa5f --- /dev/null +++ b/catch/ABM/AddKernels/CMakeLists.txt @@ -0,0 +1,8 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + add.cc +) + +hip_add_exe_to_target(NAME ABMAddKernels + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/catch/ABM/AddKernels/add.cc b/catch/ABM/AddKernels/add.cc new file mode 100644 index 000000000..1b7c56cdf --- /dev/null +++ b/catch/ABM/AddKernels/add.cc @@ -0,0 +1,41 @@ +#include +#include + +template __global__ void add(T* a, T* b, T* c, size_t size) { + size_t i = threadIdx.x; + if (i < size) c[i] = a[i] + b[i]; +} + +TEMPLATE_TEST_CASE("ABM_AddKernel_MultiTypeMultiSize", "", int, long, float, long long, double) { + auto size = GENERATE(as{}, 100, 500, 1000); + TestType *d_a, *d_b, *d_c; + auto res = hipMalloc(&d_a, sizeof(TestType) * size); + REQUIRE(res == hipSuccess); + res = hipMalloc(&d_b, sizeof(TestType) * size); + REQUIRE(res == hipSuccess); + res = hipMalloc(&d_c, sizeof(TestType) * size); + REQUIRE(res == hipSuccess); + + std::vector a, b, c; + for (size_t i = 0; i < size; i++) { + a.push_back(i + 1); + b.push_back(i + 1); + c.push_back(2 * (i + 1)); + } + + res = hipMemcpy(d_a, a.data(), sizeof(TestType) * size, hipMemcpyHostToDevice); + REQUIRE(res == hipSuccess); + res = hipMemcpy(d_b, b.data(), sizeof(TestType) * size, hipMemcpyHostToDevice); + REQUIRE(res == hipSuccess); + + hipLaunchKernelGGL(add, 1, size, 0, 0, d_a, d_b, d_c, size); + HIP_CHECK(hipGetLastError()); + + res = hipMemcpy(a.data(), d_c, sizeof(TestType) * size, hipMemcpyDeviceToHost); + REQUIRE(res == hipSuccess); + + HIP_CHECK(hipFree(d_a)); + HIP_CHECK(hipFree(d_b)); + HIP_CHECK(hipFree(d_c)); + REQUIRE(a == c); +} diff --git a/catch/ABM/CMakeLists.txt b/catch/ABM/CMakeLists.txt new file mode 100644 index 000000000..53fa9e8b0 --- /dev/null +++ b/catch/ABM/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(AddKernels) diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt new file mode 100644 index 000000000..e2924a561 --- /dev/null +++ b/catch/CMakeLists.txt @@ -0,0 +1,249 @@ +cmake_minimum_required(VERSION 3.16.8) + +# to skip the simple compiler test +set(CMAKE_C_COMPILER_WORKS 1) +set(CMAKE_CXX_COMPILER_WORKS 1) + +project(hiptests) + + +# Check if platform and compiler are set +if(HIP_PLATFORM STREQUAL "amd") + if(HIP_COMPILER STREQUAL "nvcc") + message(FATAL_ERROR "Unexpected HIP_COMPILER:${HIP_COMPILER} is set for HIP_PLATFOR:amd") + endif() +elseif(HIP_PLATFORM STREQUAL "nvidia") + if(DEFINED HIP_COMPILER AND NOT HIP_COMPILER STREQUAL "nvcc") + message(FATAL_ERROR "Unexpected HIP_COMPILER: ${HIP_COMPILER} is set for HIP_PLATFORM:nvidia") + endif() +else() + message(FATAL_ERROR "Unexpected HIP_PLATFORM: " ${HIP_PLATFORM}) +endif() + +# Set HIP Path +if(NOT DEFINED HIP_PATH) + if(DEFINED ENV{HIP_PATH}) + set(HIP_PATH $ENV{HIP_PATH} CACHE STRING "HIP Path") + else() + set(HIP_PATH "${PROJECT_BINARY_DIR}") + endif() +endif() +message(STATUS "HIP Path: ${HIP_PATH}") + +# Set ROCM Path +if(NOT DEFINED ROCM_PATH) + if(DEFINED ENV{ROCM_PATH}) + set(ROCM_PATH $ENV{ROCM_PATH} CACHE STRING "ROCM Path") + else() + cmake_path(GET HIP_PATH PARENT_PATH ROCM_PATH) + if (NOT EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator") + set(ROCM_PATH "/opt/rocm/") + endif() + endif() +endif() +file(TO_CMAKE_PATH "${ROCM_PATH}" ROCM_PATH) +message(STATUS "ROCM Path: ${ROCM_PATH}") + + +if(UNIX) + set(CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc") + set(CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc") + set(HIPCONFIG_EXECUTABLE "${HIP_PATH}/bin/hipconfig") + execute_process(COMMAND perl ${HIPCONFIG_EXECUTABLE} --version + OUTPUT_VARIABLE HIP_VERSION + OUTPUT_STRIP_TRAILING_WHITESPACE) +else() + # using cmake_path as it handles path correctly. + # Set both compilers else windows cmake complains of mismatch + cmake_path(SET CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc.bat") + cmake_path(SET CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc.bat") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${ROCM_PATH}") + set(HIPCONFIG_EXECUTABLE "${HIP_PATH}/bin/hipconfig.bat") + execute_process(COMMAND ${HIPCONFIG_EXECUTABLE} --version + OUTPUT_VARIABLE HIP_VERSION + OUTPUT_STRIP_TRAILING_WHITESPACE) +endif() +# enforce c++17 for all tests +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++17") + +string(REPLACE "." ";" VERSION_LIST ${HIP_VERSION}) +list(GET VERSION_LIST 0 HIP_VERSION_MAJOR) +list(GET VERSION_LIST 1 HIP_VERSION_MINOR) +list(GET VERSION_LIST 2 HIP_VERSION_PATCH_GITHASH) +string(REPLACE "-" ";" VERSION_LIST ${HIP_VERSION_PATCH_GITHASH}) +list(GET VERSION_LIST 0 HIP_VERSION_PATCH) + +if(NOT DEFINED CATCH2_PATH) + if(DEFINED ENV{CATCH2_PATH}) + set(CATCH2_PATH $ENV{CATCH2_PATH} CACHE STRING "Catch2 Path") + else() + set(CATCH2_PATH "${CMAKE_CURRENT_LIST_DIR}/external/Catch2") + endif() +endif() +message(STATUS "Catch2 Path: ${CATCH2_PATH}") + +# Set JSON Parser path +if(NOT DEFINED JSON_PARSER) + if(DEFINED ENV{JSON_PARSER}) + set(JSON_PARSER $ENV{JSON_PARSER} CACHE STRING "JSON Parser Path") + else() + set(JSON_PARSER "${CMAKE_CURRENT_LIST_DIR}/external/picojson") + endif() +endif() + +message(STATUS "Searching Catch2 in: ${CMAKE_CURRENT_LIST_DIR}/external") +find_package(Catch2 REQUIRED + PATHS + ${CMAKE_CURRENT_LIST_DIR}/external + PATH_SUFFIXES + Catch2/cmake/Catch2 +) +include(Catch) +include(CTest) + +# path used for generating the *_include.cmake file +set(CATCH2_INCLUDE ${CATCH2_PATH}/cmake/Catch2/catch_include.cmake.in) + +include_directories( + ${CATCH2_PATH} + "./include" + "./kernels" + ${HIP_PATH}/include + ${JSON_PARSER} +) + +option(RTC_TESTING "Run tests using HIP RTC to compile the kernels" OFF) +if (RTC_TESTING) + add_definitions(-DRTC_TESTING=ON) +endif() +add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") + +set(CATCH_BUILD_DIR catch_tests) +file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/hipTestMain) +file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake + DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script) +set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script/CatchAddTests.cmake) + +if (WIN32) + configure_file(catchProp_in_rc.in ${CMAKE_CURRENT_BINARY_DIR}/catchProp.rc @ONLY) + cmake_path(SET LLVM_RC_PATH "${HIP_PATH}/../lc/bin/llvm-rc.exe") + cmake_path(SET LLVM_RC_PATH NORMALIZE "${LLVM_RC_PATH}") + + # generates the .res files to be used by executables to populate the properties + # expects LC folder with clang, llvm-rc to be present one level up of HIP + execute_process(COMMAND ${LLVM_RC_PATH} ${CMAKE_CURRENT_BINARY_DIR}/catchProp.rc + OUTPUT_VARIABLE RC_OUTPUT) + set(PROP_RC ${CMAKE_CURRENT_BINARY_DIR}) +endif() + +if(HIP_PLATFORM MATCHES "amd" AND HIP_COMPILER MATCHES "clang") + add_compile_options(-Wall -Wextra -pedantic -Werror -Wno-deprecated) +endif() + +cmake_policy(PUSH) +if(POLICY CMP0037) + cmake_policy(SET CMP0037 OLD) +endif() + +# Turn off CMAKE_HIP_ARCHITECTURES Feature if cmake version is 3.21+ +if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.21.0) + set(CMAKE_HIP_ARCHITECTURES OFF) +endif() +message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") + +# Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" +# having space at the start/end of OFFLOAD_ARCH_STR can cause build failures +# Identify the GPU Targets. +# This is done due to limitation of rocm_agent_enumerator +# While building test parallelly, rocm_agent_enumerator can fail and give out an empty target +# That results in hipcc building the test for gfx803 (the default target) +# preference to pass arch - +# OFFLOAD_ARCH_STR +# ENV{HCC_AMDGPU_TARGET} +# rocm_agent_enumerator +if(NOT DEFINED OFFLOAD_ARCH_STR + AND NOT DEFINED ENV{HCC_AMDGPU_TARGET} + AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" + AND HIP_PLATFORM STREQUAL "amd" AND UNIX) + execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator + OUTPUT_VARIABLE HIP_GPU_ARCH + RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT) + # Trim out gfx000 + string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH}) + if (NOT HIP_GPU_ARCH STREQUAL "") + string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN) + # If string has more gfx target except gfx000 + if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1) + string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}") + set(OFFLOAD_ARCH_STR "") + foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST}) + set(OFFLOAD_ARCH_STR "--offload-arch=${_hip_gpu_arch} ${OFFLOAD_ARCH_STR}") + endforeach() + endif() + else() + message(STATUS "ROCm Agent Enumurator found no valid architectures") + endif() +elseif(DEFINED OFFLOAD_ARCH_STR) + string(REPLACE "--offload-arch=" "" HIP_GPU_ARCH_LIST ${OFFLOAD_ARCH_STR}) +endif() + +if(DEFINED OFFLOAD_ARCH_STR) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ") +elseif(DEFINED ENV{HCC_AMDGPU_TARGET}) + # hipcc pl script appends it to the options + set(OFFLOAD_ARCH_STR " --offload-arch=$ENV{HCC_AMDGPU_TARGET}") + set(HIP_GPU_ARCH_LIST $ENV{HCC_AMDGPU_TARGET}) +endif() +message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}") + +# prints the catch info to a file +string(TIMESTAMP _timestamp UTC) +set(_catchInfo "# Auto-generated by cmake on ${_timestamp} UTC\n") +set(_catchInfo ${_catchInfo} "HIP_VERSION=${HIP_VERSION}\n") +set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n") +set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n") +file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo}) + +# Enable device lambda on nvidia platforms +if(HIP_COMPILER MATCHES "nvcc") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --extended-lambda") +endif() + +# Disable CXX extensions (gnu++11 etc) +set(CMAKE_CXX_EXTENSIONS OFF) + +add_custom_target(build_tests) + + +# Tests folder +add_subdirectory(unit ${CATCH_BUILD_DIR}/unit) +add_subdirectory(ABM ${CATCH_BUILD_DIR}/ABM) +add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels) +add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain) +add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) +add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) +if(UNIX) + add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) +endif() + +cmake_policy(POP) + +# packaging the tests +# make package_test to generate packages for test +set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/) +configure_file(packaging/hip-tests.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +if(UNIX) +add_custom_target(package_test COMMAND ${CMAKE_COMMAND} . + COMMAND rm -rf *.deb *.rpm *.tar.gz + COMMAND make package + COMMAND cp *.deb ${PROJECT_BINARY_DIR} + COMMAND cp *.rpm ${PROJECT_BINARY_DIR} + COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} + WORKING_DIRECTORY ${BUILD_DIR}) +else() +file(TO_NATIVE_PATH ${PROJECT_BINARY_DIR} CATCH_BINARY_DIR) +add_custom_target(package_test COMMAND ${CMAKE_COMMAND} . + COMMAND cpack + COMMAND copy *.zip ${CATCH_BINARY_DIR} + WORKING_DIRECTORY ${BUILD_DIR}) +endif() diff --git a/catch/README.md b/catch/README.md new file mode 100644 index 000000000..5688f558d --- /dev/null +++ b/catch/README.md @@ -0,0 +1,204 @@ +# HIP Tests - with Catch2 + +## Intro and Motivation +HIP Tests were using HIT framework (a custom framework tailored for HIP) to add, build and run tests. As time progressed the frame got big and took substantial amount of effort to maintain and extend. It also took substantial amount of time to configure. We took this oppurtunity to rewrite the HIP's testing framework and porting the test infra to Catch2 format. + +## How to write tests +Tests in Catch2 are declared via ```TEST_CASE```. + +[Please read the Catch2 documentation on how to write test cases](https://github.com/catchorg/Catch2/blob/v2.13.6/docs/tutorial.md#top) + +[Catch2 Detailed Reference](https://github.com/catchorg/Catch2/blob/v2.13.6/docs/Readme.md#top) + +## Taking care of existing features +- Don’t build on platform: EXCLUDE_(HIP_PLATFORM/HIP_RUNTIME), can be done via CMAKE. Adding source in if(HIP_PLATFORM == amd/nvidia). +- HIPCC_OPTIONS/CLANG Options: Can be done via: set_source_files_properties(src.cc PROPERTIES COMPILE_FLAGS “…”). +- Additional libraries: Can be done via target_link_libraries() +- Multiple runs with different args: This can be done by Catch’s Feature: GENERATE(…) +Running Subtest: ctest –R “...” (Regex to match the subtest name) + +## New Features +- Skip test without recompiling tests, by addition of a json file. Default name is ```config.json``` , this can be overridden by using the variable ```HIP_CATCH_EXCLUDE_FILE=some_config.json```. +- Json file supports regex. Ex: All tests which has the word ‘Memset’ can be skipped using ‘*Memset*’ +- Support multiple skip test list which can be set via environment variable, so you can have multiple files containing different skip test lists and can pick and choose among them depending on your platform and os. +- Better CI integration via xunit compatible output + +## Testing Context +HIP testing framework gives you a context for each test. This context will have useful information about the environment your test is running. + +Some useful functions are: +- `bool isWindows()` : true if os is windows +- `bool isLinux()` : true if os is linux +- `bool isAmd()` : true if platform is AMD +- `bool isNvidia()` : true if platform is NVIDIA + +This information can be accessed in any test via using: `TestContext::get().isAmd()`. + +## Adding test for a specific platform +There might be some functionality which is not present on some platforms. Those tests can be hidden inside following macros. + +- ```HT_AMD``` is 1 when tests are running on AMD platform and 0 on NVIDIA. +- ```HT_NVIDIA``` is 1 when tests are running on NVIDIA platform and 0 on AMD + +Usage: + +```cpp +#if HT_AMD +TEST_CASE("hipExtAPIs") { + // ... +} +#endif +``` + +## Config file schema +Some tests can be skipped using a config file placed in hipTestMain/config folder. Multiple config files can be defined for different configurations. +The naming convention for the file needs to be "config_platform_os_archname.json" +Platform and os are mandatory. +Arch name is optional and takes precedence while loading the json file. +Currently the json files need to be manually chosen by the executor for the architecture of choice. + +example: +config_amd_windows.json +config_nvidia_windows.json + +The schema of the json file is as follows: +```json +{ + "DisabledTests": + [ + "TestName1", + "TestName2", + ... + ] +} +``` + +## Environment Variables +- `HIP_CATCH_EXCLUDE_FILE` : This variable can be set to the config file name or full path. Disabled tests will be read from this. +- `HT_LOG_ENABLE` : This is for debugging the HIP Test Framework itself. Setting it to 1, all `LogPrintf` will be printed on screen + +## Test Macros +### Single Thread Macros +These macros are to be used when your test is calling HIP APIs via the main thread. + +- `HIP_CHECK` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. + + - Usage: ```HIP_CHECK(hipMalloc(&dPtr, 10));``` + +- ```HIP_CHECK_ERROR``` : This macro takes in a HIP API and tests its result against a provided result. This can be used when the API is expected to fail with a particular result. + + - Usage: ```HIP_CHECK_ERROR(hipMalloc(&dPtr, 0), hipErrorInvalidValue);``` + +- ```HIPRTC_CHECK``` : This macro takes in a HIPRTC API and tests its result against HIPRTC_SUCCESS. + + - Usage: ```HIPRTC_CHECK(hiprtcCompileProgram(prog, count, options));``` + +- ```HIP_ASSERT``` : This macro takes in a bool condition as input and does a ```REQUIRE``` on the condition. + + - Usage: ```HIP_ASSERT(result == 10);``` + +### Multi Thread Macros +These macros are to be used when you call HIP APIs in a multi threaded way. They exist because Catch2 ```REQUIRE``` and ```CHECK``` macros can not handle multi threaded calls. To solve this problem, two macros are added```HIP_CHECK_THREAD``` and ```REQUIRE_THREAD``` which can be used to check result of HIP APIs and test assertions respectively. The results can be validate after the threads join via ```HIP_CHECK_THREAD_FINALIZE```. + +Note: These should used in ```std::thread``` only. For multi proc guidelines look at [MultiProc Macros](#multi-process-macros) and [SpawnProc Class](#multiproc-management-class) + +- ```HIP_CHECK_THREAD``` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. It can also tell other threads if an error has occured in one of the HIP API and can prematurely stop the threads. + +- ```REQUIRE_THREAD``` : This macro takes in a bool condition and tests for its result to be true. If this check fails, it can signal other threads to terminate early. + +- ```HIP_CHECK_THREAD_FINALIZE``` : This macro checks for the results logged by ```HIP_CHECK_THREAD```. This needs to be called after the threads have joined. + +Please also note that you can not return values in functions calling ```HIP_CHECK_THREAD``` or ```REQUIRE_THREAD``` macro. + + Usage: + + ```cpp + auto threadFunc = []() { + int *dPtr{nullptr}; + HIP_CHECK_THREAD(hipMalloc(&dPtr, 10)); + REQUIRE_THREAD(dPtr != nullptr); + // Some other work + }; + + // Launch threads + std::vector threadPool; + for(...) { + threadPool.emplace_back(std::thread(threadFunc)); + } + + // Join threads + for(auto &i : threadPool) { + i.join(); + } + + // Validate all results + HIP_CHECK_THREAD_FINALIZE(); + ``` + +### Skipping Tests if certain criteria is not met +If there arises a condition where certain flag is disabled and due to which a test can not run at that time, the following macro can be of use. It will highlight the test in ctest report as well. + +- ```HIP_SKIP_TEST``` : The api takes in an input of the reason as well and prints out the line HIP_SKIP_THIS_TEST. This causes ctest to mark the test as skipped and the test shows up in the report as skipped prompting proper response from the team. + + Usage: + + ```cpp + TEST_CASE("TestOnlyOnXnack") { + if(!XNACKEnabled) { + HipTest::HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); + return; + } + // Rest of test functionality + } + ``` + +### Multi Process Macros +These macros are to be called in multi process tests, inside a process which gets spawned. The reasoning is the same, Catch2 does not support multi process checks. + +- ```HIPCHECK``` : Same as ```HIP_CHECK``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +- ```HIPASSERT``` : Same as ```HIP_ASSERT``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +## MultiProc Management Class +There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output. +Sample Usage: +```cpp +hip::SpawnProc proc(, ); +REQUIRE(0 == proc.run()); // Test of return value of the proc +REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process +``` +The process must be a standalone exe inside the same folder as other tests. + +## Enabling New Tests +Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default. + +## Building a single test +```bash +hipcc -I/tests/catch/include /tests/catch/hipTestMain/standalone_main.cc -I/tests/catch/external/Catch2 -g -o +``` + +## Debugging support +Catch2 allows multiple ways in which you can debug the test case. +- `-b` options breaks into a debugger as soon as there is a failure encountered [Catch2 Options Reference](https://github.com/catchorg/Catch2/blob/devel/docs/command-line.md#breaking-into-the-debugger) +- Catch2 provided [logging macro](https://github.com/catchorg/Catch2/blob/v2.13.6/docs/logging.md#top) that print useful information on test case failure +- User can also call [CATCH_BREAK_INTO_DEBUGGER](https://github.com/catchorg/Catch2/blob/devel/docs/configuration.md#overriding-catchs-debug-break--b) macro to break at a certain point in a test case. +- User can also mention filename.cc:__LineNumber__ to break into a test case via gdb. + +## External Libs being used +- [Catch2](https://github.com/catchorg/Catch2) - Testing framework +- [picojson](https://github.com/kazuho/picojson) - For config file parsing + +# Testing Guidelines +Tests fall in 5 categories and its file name prefix are as follows: + - Unit tests (Prefix: Unit_\*API\*_\*Optional Scenario\*, example : Unit_hipMalloc_Negative or Unit_hipMalloc): Unit Tests are simplest test for an API, the target here is to test the API with different types of input and different ways of calling. + - Application Behavior Modelling tests (Prefix: ABM_\*Intent\*_\*Optional Scenario\*, example: ABM_ModuleLoadAndRun): ABM tests are used to model a specific use case of HIP APIs, either seen in a customer app or a general purpose app. It mimics the calling behavior seen in aforementioned app. + - Stress/Scale tests (Prefix: Stress_\*API\*_\*Intent\*_\*Optional Scenario\*, example: Stress_hipMemset_ExhaustVRAM): These tests are used to see the behavior of HIP APIs in edge scenarios, for example what happens when we have exhausted vram and do a hipMalloc or run many instances of same API in parallel. + - Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment + - Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc y): Performance tests are used to get results of HIP APIs. + +# General Guidelines: + - Do not use the catch2 tags. Tags wont be used for filtering + - Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails) + - Check return of each HIP API and fail whenever there is a misma tch with hipSuccess or hiprtcSuccess. + - Each Category of test will hav e its own exe and catch_discover_test macro will be called on it to discover its tests + - Optional Scenario in test names are optional. For example you can test all Scenarios of hipMalloc API in one file, you can name the file Unit_hipMalloc, if you are having a file just for negative scenarios you can name it as Unit_hipMalloc_Negative. diff --git a/catch/TypeQualifiers/CMakeLists.txt b/catch/TypeQualifiers/CMakeLists.txt new file mode 100644 index 000000000..0219c0021 --- /dev/null +++ b/catch/TypeQualifiers/CMakeLists.txt @@ -0,0 +1,8 @@ +# Common Tests +set(TEST_SRC + hipManagedKeyword.cc +) + +hip_add_exe_to_target(NAME TypeQualifiers + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/catch/TypeQualifiers/hipManagedKeyword.cc b/catch/TypeQualifiers/hipManagedKeyword.cc new file mode 100644 index 000000000..4d6a97b11 --- /dev/null +++ b/catch/TypeQualifiers/hipManagedKeyword.cc @@ -0,0 +1,76 @@ +/* + Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +/* + This testcase verifies the hipManagedKeyword basic scenario + */ + +#include +#include + +#define N 1048576 +__managed__ float A[N]; // Accessible by ALL CPU and GPU functions !!! +__managed__ float B[N]; +__managed__ int x = 0; + +__global__ void add(const float *A, float *B) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < N; i += stride) + B[i] = A[i] + B[i]; +} + +__global__ void GPU_func() { + x++; +} + +TEST_CASE("Unit_hipManagedKeyword_SingleGpu") { + for (int i = 0; i < N; i++) { + A[i] = 1.0f; + B[i] = 2.0f; + } + + int blockSize = 256; + int numBlocks = (N + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, static_cast(A), + static_cast(B)); + + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + + float maxError = 0.0f; + for (int i = 0; i < N; i++) + maxError = fmax(maxError, fabs(B[i]-3.0f)); + + REQUIRE(maxError == 0.0f); +} + +TEST_CASE("Unit_hipManagedKeyword_MultiGpu") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipSetDevice(i)); + GPU_func<<< 1, 1 >>>(); + HIP_CHECK(hipDeviceSynchronize()); + } + REQUIRE(x == numDevices); +} diff --git a/catch/catchProp_in_rc.in b/catch/catchProp_in_rc.in new file mode 100644 index 000000000..74d202818 --- /dev/null +++ b/catch/catchProp_in_rc.in @@ -0,0 +1,40 @@ +#include + +#define HIP_VERSION "@HIP_VERSION@" +#define HIP_VERSION_MAJOR @HIP_VERSION_MAJOR@ +#define HIP_VERSION_MINOR @HIP_VERSION_MINOR@ +#define HIP_VERSION_PATCH @HIP_VERSION_PATCH@ + +VS_VERSION_INFO VERSIONINFO +FILEVERSION HIP_VERSION_MAJOR, HIP_VERSION_MINOR , HIP_VERSION_PATCH +PRODUCTVERSION 10,1 +FILEFLAGSMASK 0x3fL +#ifdef _DEBUG + FILEFLAGS VS_FF_DEBUG +#else + FILEFLAGS 0x0L +#endif +FILEOS VOS_NT_WINDOWS32 +FILETYPE VFT_APP +FILESUBTYPE VFT2_UNKNOWN +BEGIN + BLOCK "StringFileInfo" + BEGIN + BLOCK "040904b0" + BEGIN + VALUE "CompanyName", "Advanced Micro Devices Inc.\0" + VALUE "FileDescription", "HIP unit tests" + VALUE "FileVersion", "amdhip64.dll" HIP_VERSION + VALUE "LegalCopyright", "Copyright (C) 2022 Advanced Micro Devices Inc.\0" + VALUE "ProductName", "HIP unit tests" + VALUE "ProductVersion", HIP_VERSION + VALUE "Comments", "\0" + VALUE "InternalName", "HIP unit tests" + END + END + BLOCK "VarFileInfo" + BEGIN + VALUE "Translation", 0x0409, 1200 + END +END +/* End of Version info */ diff --git a/catch/external/Catch2/LICENSE.txt b/catch/external/Catch2/LICENSE.txt new file mode 100644 index 000000000..36b7cd93c --- /dev/null +++ b/catch/external/Catch2/LICENSE.txt @@ -0,0 +1,23 @@ +Boost Software License - Version 1.0 - August 17th, 2003 + +Permission is hereby granted, free of charge, to any person or organization +obtaining a copy of the software and accompanying documentation covered by +this license (the "Software") to use, reproduce, display, distribute, +execute, and transmit the Software, and to prepare derivative works of the +Software, and to permit third-parties to whom the Software is furnished to +do so, all subject to the following: + +The copyright notices in the Software and this entire statement, including +the above license grant, this restriction and the following disclaimer, +must be included in all copies of the Software, in whole or in part, and +all derivative works of the Software, unless such copies or derivative +works are solely in the form of machine-executable object code generated by +a source language processor. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE, TITLE AND NON-INFRINGEMENT. IN NO EVENT +SHALL THE COPYRIGHT HOLDERS OR ANYONE DISTRIBUTING THE SOFTWARE BE LIABLE +FOR ANY DAMAGES OR OTHER LIABILITY, WHETHER IN CONTRACT, TORT OR OTHERWISE, +ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +DEALINGS IN THE SOFTWARE. diff --git a/catch/external/Catch2/catch.hpp b/catch/external/Catch2/catch.hpp new file mode 100644 index 000000000..6897dae0d --- /dev/null +++ b/catch/external/Catch2/catch.hpp @@ -0,0 +1,17881 @@ +/* + * Catch v2.13.4 + * Generated: 2020-12-29 14:48:00.116107 + * ---------------------------------------------------------- + * This file has been merged from multiple headers. Please don't edit it directly + * Copyright (c) 2020 Two Blue Cubes Ltd. All rights reserved. + * + * Distributed under the Boost Software License, Version 1.0. (See accompanying + * file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + */ +#ifndef TWOBLUECUBES_SINGLE_INCLUDE_CATCH_HPP_INCLUDED +#define TWOBLUECUBES_SINGLE_INCLUDE_CATCH_HPP_INCLUDED +// start catch.hpp + + +#define CATCH_VERSION_MAJOR 2 +#define CATCH_VERSION_MINOR 13 +#define CATCH_VERSION_PATCH 4 + +#ifdef __clang__ +# pragma clang system_header +#elif defined __GNUC__ +# pragma GCC system_header +#endif + +// start catch_suppress_warnings.h + +#ifdef __clang__ +# ifdef __ICC // icpc defines the __clang__ macro +# pragma warning(push) +# pragma warning(disable: 161 1682) +# else // __ICC +# pragma clang diagnostic push +# pragma clang diagnostic ignored "-Wpadded" +# pragma clang diagnostic ignored "-Wswitch-enum" +# pragma clang diagnostic ignored "-Wcovered-switch-default" +# endif +#elif defined __GNUC__ + // Because REQUIREs trigger GCC's -Wparentheses, and because still + // supported version of g++ have only buggy support for _Pragmas, + // Wparentheses have to be suppressed globally. +# pragma GCC diagnostic ignored "-Wparentheses" // See #674 for details + +# pragma GCC diagnostic push +# pragma GCC diagnostic ignored "-Wunused-variable" +# pragma GCC diagnostic ignored "-Wpadded" +#endif +// end catch_suppress_warnings.h +#if defined(CATCH_CONFIG_MAIN) || defined(CATCH_CONFIG_RUNNER) +# define CATCH_IMPL +# define CATCH_CONFIG_ALL_PARTS +#endif + +// In the impl file, we want to have access to all parts of the headers +// Can also be used to sanely support PCHs +#if defined(CATCH_CONFIG_ALL_PARTS) +# define CATCH_CONFIG_EXTERNAL_INTERFACES +# if defined(CATCH_CONFIG_DISABLE_MATCHERS) +# undef CATCH_CONFIG_DISABLE_MATCHERS +# endif +# if !defined(CATCH_CONFIG_ENABLE_CHRONO_STRINGMAKER) +# define CATCH_CONFIG_ENABLE_CHRONO_STRINGMAKER +# endif +#endif + +#if !defined(CATCH_CONFIG_IMPL_ONLY) +// start catch_platform.h + +#ifdef __APPLE__ +# include +# if TARGET_OS_OSX == 1 +# define CATCH_PLATFORM_MAC +# elif TARGET_OS_IPHONE == 1 +# define CATCH_PLATFORM_IPHONE +# endif + +#elif defined(linux) || defined(__linux) || defined(__linux__) +# define CATCH_PLATFORM_LINUX + +#elif defined(WIN32) || defined(__WIN32__) || defined(_WIN32) || defined(_MSC_VER) || defined(__MINGW32__) +# define CATCH_PLATFORM_WINDOWS +#endif + +// end catch_platform.h + +#ifdef CATCH_IMPL +# ifndef CLARA_CONFIG_MAIN +# define CLARA_CONFIG_MAIN_NOT_DEFINED +# define CLARA_CONFIG_MAIN +# endif +#endif + +// start catch_user_interfaces.h + +namespace Catch { + unsigned int rngSeed(); +} + +// end catch_user_interfaces.h +// start catch_tag_alias_autoregistrar.h + +// start catch_common.h + +// start catch_compiler_capabilities.h + +// Detect a number of compiler features - by compiler +// The following features are defined: +// +// CATCH_CONFIG_COUNTER : is the __COUNTER__ macro supported? +// CATCH_CONFIG_WINDOWS_SEH : is Windows SEH supported? +// CATCH_CONFIG_POSIX_SIGNALS : are POSIX signals supported? +// CATCH_CONFIG_DISABLE_EXCEPTIONS : Are exceptions enabled? +// **************** +// Note to maintainers: if new toggles are added please document them +// in configuration.md, too +// **************** + +// In general each macro has a _NO_ form +// (e.g. CATCH_CONFIG_NO_POSIX_SIGNALS) which disables the feature. +// Many features, at point of detection, define an _INTERNAL_ macro, so they +// can be combined, en-mass, with the _NO_ forms later. + +#ifdef __cplusplus + +# if (__cplusplus >= 201402L) || (defined(_MSVC_LANG) && _MSVC_LANG >= 201402L) +# define CATCH_CPP14_OR_GREATER +# endif + +# if (__cplusplus >= 201703L) || (defined(_MSVC_LANG) && _MSVC_LANG >= 201703L) +# define CATCH_CPP17_OR_GREATER +# endif + +#endif + +// We have to avoid both ICC and Clang, because they try to mask themselves +// as gcc, and we want only GCC in this block +#if defined(__GNUC__) && !defined(__clang__) && !defined(__ICC) && !defined(__CUDACC__) +# define CATCH_INTERNAL_START_WARNINGS_SUPPRESSION _Pragma( "GCC diagnostic push" ) +# define CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION _Pragma( "GCC diagnostic pop" ) + +# define CATCH_INTERNAL_IGNORE_BUT_WARN(...) (void)__builtin_constant_p(__VA_ARGS__) + +#endif + +#if defined(__clang__) + +# define CATCH_INTERNAL_START_WARNINGS_SUPPRESSION _Pragma( "clang diagnostic push" ) +# define CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION _Pragma( "clang diagnostic pop" ) + +// As of this writing, IBM XL's implementation of __builtin_constant_p has a bug +// which results in calls to destructors being emitted for each temporary, +// without a matching initialization. In practice, this can result in something +// like `std::string::~string` being called on an uninitialized value. +// +// For example, this code will likely segfault under IBM XL: +// ``` +// REQUIRE(std::string("12") + "34" == "1234") +// ``` +// +// Therefore, `CATCH_INTERNAL_IGNORE_BUT_WARN` is not implemented. +# if !defined(__ibmxl__) && !defined(__CUDACC__) +# define CATCH_INTERNAL_IGNORE_BUT_WARN(...) (void)__builtin_constant_p(__VA_ARGS__) /* NOLINT(cppcoreguidelines-pro-type-vararg, hicpp-vararg) */ +# endif + +# define CATCH_INTERNAL_SUPPRESS_GLOBALS_WARNINGS \ + _Pragma( "clang diagnostic ignored \"-Wexit-time-destructors\"" ) \ + _Pragma( "clang diagnostic ignored \"-Wglobal-constructors\"") + +# define CATCH_INTERNAL_SUPPRESS_PARENTHESES_WARNINGS \ + _Pragma( "clang diagnostic ignored \"-Wparentheses\"" ) + +# define CATCH_INTERNAL_SUPPRESS_UNUSED_WARNINGS \ + _Pragma( "clang diagnostic ignored \"-Wunused-variable\"" ) + +# define CATCH_INTERNAL_SUPPRESS_ZERO_VARIADIC_WARNINGS \ + _Pragma( "clang diagnostic ignored \"-Wgnu-zero-variadic-macro-arguments\"" ) + +# define CATCH_INTERNAL_SUPPRESS_UNUSED_TEMPLATE_WARNINGS \ + _Pragma( "clang diagnostic ignored \"-Wunused-template\"" ) + +#endif // __clang__ + +//////////////////////////////////////////////////////////////////////////////// +// Assume that non-Windows platforms support posix signals by default +#if !defined(CATCH_PLATFORM_WINDOWS) + #define CATCH_INTERNAL_CONFIG_POSIX_SIGNALS +#endif + +//////////////////////////////////////////////////////////////////////////////// +// We know some environments not to support full POSIX signals +#if defined(__CYGWIN__) || defined(__QNX__) || defined(__EMSCRIPTEN__) || defined(__DJGPP__) + #define CATCH_INTERNAL_CONFIG_NO_POSIX_SIGNALS +#endif + +#ifdef __OS400__ +# define CATCH_INTERNAL_CONFIG_NO_POSIX_SIGNALS +# define CATCH_CONFIG_COLOUR_NONE +#endif + +//////////////////////////////////////////////////////////////////////////////// +// Android somehow still does not support std::to_string +#if defined(__ANDROID__) +# define CATCH_INTERNAL_CONFIG_NO_CPP11_TO_STRING +# define CATCH_INTERNAL_CONFIG_ANDROID_LOGWRITE +#endif + +//////////////////////////////////////////////////////////////////////////////// +// Not all Windows environments support SEH properly +#if defined(__MINGW32__) +# define CATCH_INTERNAL_CONFIG_NO_WINDOWS_SEH +#endif + +//////////////////////////////////////////////////////////////////////////////// +// PS4 +#if defined(__ORBIS__) +# define CATCH_INTERNAL_CONFIG_NO_NEW_CAPTURE +#endif + +//////////////////////////////////////////////////////////////////////////////// +// Cygwin +#ifdef __CYGWIN__ + +// Required for some versions of Cygwin to declare gettimeofday +// see: http://stackoverflow.com/questions/36901803/gettimeofday-not-declared-in-this-scope-cygwin +# define _BSD_SOURCE +// some versions of cygwin (most) do not support std::to_string. Use the libstd check. +// https://gcc.gnu.org/onlinedocs/gcc-4.8.2/libstdc++/api/a01053_source.html line 2812-2813 +# if !((__cplusplus >= 201103L) && defined(_GLIBCXX_USE_C99) \ + && !defined(_GLIBCXX_HAVE_BROKEN_VSWPRINTF)) + +# define CATCH_INTERNAL_CONFIG_NO_CPP11_TO_STRING + +# endif +#endif // __CYGWIN__ + +//////////////////////////////////////////////////////////////////////////////// +// Visual C++ +#if defined(_MSC_VER) + +# define CATCH_INTERNAL_START_WARNINGS_SUPPRESSION __pragma( warning(push) ) +# define CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION __pragma( warning(pop) ) + +// Universal Windows platform does not support SEH +// Or console colours (or console at all...) +# if defined(WINAPI_FAMILY) && (WINAPI_FAMILY == WINAPI_FAMILY_APP) +# define CATCH_CONFIG_COLOUR_NONE +# else +# define CATCH_INTERNAL_CONFIG_WINDOWS_SEH +# endif + +// MSVC traditional preprocessor needs some workaround for __VA_ARGS__ +// _MSVC_TRADITIONAL == 0 means new conformant preprocessor +// _MSVC_TRADITIONAL == 1 means old traditional non-conformant preprocessor +# if !defined(__clang__) // Handle Clang masquerading for msvc +# if !defined(_MSVC_TRADITIONAL) || (defined(_MSVC_TRADITIONAL) && _MSVC_TRADITIONAL) +# define CATCH_INTERNAL_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR +# endif // MSVC_TRADITIONAL +# endif // __clang__ + +#endif // _MSC_VER + +#if defined(_REENTRANT) || defined(_MSC_VER) +// Enable async processing, as -pthread is specified or no additional linking is required +# define CATCH_INTERNAL_CONFIG_USE_ASYNC +#endif // _MSC_VER + +//////////////////////////////////////////////////////////////////////////////// +// Check if we are compiled with -fno-exceptions or equivalent +#if defined(__EXCEPTIONS) || defined(__cpp_exceptions) || defined(_CPPUNWIND) +# define CATCH_INTERNAL_CONFIG_EXCEPTIONS_ENABLED +#endif + +//////////////////////////////////////////////////////////////////////////////// +// DJGPP +#ifdef __DJGPP__ +# define CATCH_INTERNAL_CONFIG_NO_WCHAR +#endif // __DJGPP__ + +//////////////////////////////////////////////////////////////////////////////// +// Embarcadero C++Build +#if defined(__BORLANDC__) + #define CATCH_INTERNAL_CONFIG_POLYFILL_ISNAN +#endif + +//////////////////////////////////////////////////////////////////////////////// + +// Use of __COUNTER__ is suppressed during code analysis in +// CLion/AppCode 2017.2.x and former, because __COUNTER__ is not properly +// handled by it. +// Otherwise all supported compilers support COUNTER macro, +// but user still might want to turn it off +#if ( !defined(__JETBRAINS_IDE__) || __JETBRAINS_IDE__ >= 20170300L ) + #define CATCH_INTERNAL_CONFIG_COUNTER +#endif + +//////////////////////////////////////////////////////////////////////////////// + +// RTX is a special version of Windows that is real time. +// This means that it is detected as Windows, but does not provide +// the same set of capabilities as real Windows does. +#if defined(UNDER_RTSS) || defined(RTX64_BUILD) + #define CATCH_INTERNAL_CONFIG_NO_WINDOWS_SEH + #define CATCH_INTERNAL_CONFIG_NO_ASYNC + #define CATCH_CONFIG_COLOUR_NONE +#endif + +#if !defined(_GLIBCXX_USE_C99_MATH_TR1) +#define CATCH_INTERNAL_CONFIG_GLOBAL_NEXTAFTER +#endif + +// Various stdlib support checks that require __has_include +#if defined(__has_include) + // Check if string_view is available and usable + #if __has_include() && defined(CATCH_CPP17_OR_GREATER) + # define CATCH_INTERNAL_CONFIG_CPP17_STRING_VIEW + #endif + + // Check if optional is available and usable + # if __has_include() && defined(CATCH_CPP17_OR_GREATER) + # define CATCH_INTERNAL_CONFIG_CPP17_OPTIONAL + # endif // __has_include() && defined(CATCH_CPP17_OR_GREATER) + + // Check if byte is available and usable + # if __has_include() && defined(CATCH_CPP17_OR_GREATER) + # include + # if __cpp_lib_byte > 0 + # define CATCH_INTERNAL_CONFIG_CPP17_BYTE + # endif + # endif // __has_include() && defined(CATCH_CPP17_OR_GREATER) + + // Check if variant is available and usable + # if __has_include() && defined(CATCH_CPP17_OR_GREATER) + # if defined(__clang__) && (__clang_major__ < 8) + // work around clang bug with libstdc++ https://bugs.llvm.org/show_bug.cgi?id=31852 + // fix should be in clang 8, workaround in libstdc++ 8.2 + # include + # if defined(__GLIBCXX__) && defined(_GLIBCXX_RELEASE) && (_GLIBCXX_RELEASE < 9) + # define CATCH_CONFIG_NO_CPP17_VARIANT + # else + # define CATCH_INTERNAL_CONFIG_CPP17_VARIANT + # endif // defined(__GLIBCXX__) && defined(_GLIBCXX_RELEASE) && (_GLIBCXX_RELEASE < 9) + # else + # define CATCH_INTERNAL_CONFIG_CPP17_VARIANT + # endif // defined(__clang__) && (__clang_major__ < 8) + # endif // __has_include() && defined(CATCH_CPP17_OR_GREATER) +#endif // defined(__has_include) + +#if defined(CATCH_INTERNAL_CONFIG_COUNTER) && !defined(CATCH_CONFIG_NO_COUNTER) && !defined(CATCH_CONFIG_COUNTER) +# define CATCH_CONFIG_COUNTER +#endif +#if defined(CATCH_INTERNAL_CONFIG_WINDOWS_SEH) && !defined(CATCH_CONFIG_NO_WINDOWS_SEH) && !defined(CATCH_CONFIG_WINDOWS_SEH) && !defined(CATCH_INTERNAL_CONFIG_NO_WINDOWS_SEH) +# define CATCH_CONFIG_WINDOWS_SEH +#endif +// This is set by default, because we assume that unix compilers are posix-signal-compatible by default. +#if defined(CATCH_INTERNAL_CONFIG_POSIX_SIGNALS) && !defined(CATCH_INTERNAL_CONFIG_NO_POSIX_SIGNALS) && !defined(CATCH_CONFIG_NO_POSIX_SIGNALS) && !defined(CATCH_CONFIG_POSIX_SIGNALS) +# define CATCH_CONFIG_POSIX_SIGNALS +#endif +// This is set by default, because we assume that compilers with no wchar_t support are just rare exceptions. +#if !defined(CATCH_INTERNAL_CONFIG_NO_WCHAR) && !defined(CATCH_CONFIG_NO_WCHAR) && !defined(CATCH_CONFIG_WCHAR) +# define CATCH_CONFIG_WCHAR +#endif + +#if !defined(CATCH_INTERNAL_CONFIG_NO_CPP11_TO_STRING) && !defined(CATCH_CONFIG_NO_CPP11_TO_STRING) && !defined(CATCH_CONFIG_CPP11_TO_STRING) +# define CATCH_CONFIG_CPP11_TO_STRING +#endif + +#if defined(CATCH_INTERNAL_CONFIG_CPP17_OPTIONAL) && !defined(CATCH_CONFIG_NO_CPP17_OPTIONAL) && !defined(CATCH_CONFIG_CPP17_OPTIONAL) +# define CATCH_CONFIG_CPP17_OPTIONAL +#endif + +#if defined(CATCH_INTERNAL_CONFIG_CPP17_STRING_VIEW) && !defined(CATCH_CONFIG_NO_CPP17_STRING_VIEW) && !defined(CATCH_CONFIG_CPP17_STRING_VIEW) +# define CATCH_CONFIG_CPP17_STRING_VIEW +#endif + +#if defined(CATCH_INTERNAL_CONFIG_CPP17_VARIANT) && !defined(CATCH_CONFIG_NO_CPP17_VARIANT) && !defined(CATCH_CONFIG_CPP17_VARIANT) +# define CATCH_CONFIG_CPP17_VARIANT +#endif + +#if defined(CATCH_INTERNAL_CONFIG_CPP17_BYTE) && !defined(CATCH_CONFIG_NO_CPP17_BYTE) && !defined(CATCH_CONFIG_CPP17_BYTE) +# define CATCH_CONFIG_CPP17_BYTE +#endif + +#if defined(CATCH_CONFIG_EXPERIMENTAL_REDIRECT) +# define CATCH_INTERNAL_CONFIG_NEW_CAPTURE +#endif + +#if defined(CATCH_INTERNAL_CONFIG_NEW_CAPTURE) && !defined(CATCH_INTERNAL_CONFIG_NO_NEW_CAPTURE) && !defined(CATCH_CONFIG_NO_NEW_CAPTURE) && !defined(CATCH_CONFIG_NEW_CAPTURE) +# define CATCH_CONFIG_NEW_CAPTURE +#endif + +#if !defined(CATCH_INTERNAL_CONFIG_EXCEPTIONS_ENABLED) && !defined(CATCH_CONFIG_DISABLE_EXCEPTIONS) +# define CATCH_CONFIG_DISABLE_EXCEPTIONS +#endif + +#if defined(CATCH_INTERNAL_CONFIG_POLYFILL_ISNAN) && !defined(CATCH_CONFIG_NO_POLYFILL_ISNAN) && !defined(CATCH_CONFIG_POLYFILL_ISNAN) +# define CATCH_CONFIG_POLYFILL_ISNAN +#endif + +#if defined(CATCH_INTERNAL_CONFIG_USE_ASYNC) && !defined(CATCH_INTERNAL_CONFIG_NO_ASYNC) && !defined(CATCH_CONFIG_NO_USE_ASYNC) && !defined(CATCH_CONFIG_USE_ASYNC) +# define CATCH_CONFIG_USE_ASYNC +#endif + +#if defined(CATCH_INTERNAL_CONFIG_ANDROID_LOGWRITE) && !defined(CATCH_CONFIG_NO_ANDROID_LOGWRITE) && !defined(CATCH_CONFIG_ANDROID_LOGWRITE) +# define CATCH_CONFIG_ANDROID_LOGWRITE +#endif + +#if defined(CATCH_INTERNAL_CONFIG_GLOBAL_NEXTAFTER) && !defined(CATCH_CONFIG_NO_GLOBAL_NEXTAFTER) && !defined(CATCH_CONFIG_GLOBAL_NEXTAFTER) +# define CATCH_CONFIG_GLOBAL_NEXTAFTER +#endif + +// Even if we do not think the compiler has that warning, we still have +// to provide a macro that can be used by the code. +#if !defined(CATCH_INTERNAL_START_WARNINGS_SUPPRESSION) +# define CATCH_INTERNAL_START_WARNINGS_SUPPRESSION +#endif +#if !defined(CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION) +# define CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION +#endif +#if !defined(CATCH_INTERNAL_SUPPRESS_PARENTHESES_WARNINGS) +# define CATCH_INTERNAL_SUPPRESS_PARENTHESES_WARNINGS +#endif +#if !defined(CATCH_INTERNAL_SUPPRESS_GLOBALS_WARNINGS) +# define CATCH_INTERNAL_SUPPRESS_GLOBALS_WARNINGS +#endif +#if !defined(CATCH_INTERNAL_SUPPRESS_UNUSED_WARNINGS) +# define CATCH_INTERNAL_SUPPRESS_UNUSED_WARNINGS +#endif +#if !defined(CATCH_INTERNAL_SUPPRESS_ZERO_VARIADIC_WARNINGS) +# define CATCH_INTERNAL_SUPPRESS_ZERO_VARIADIC_WARNINGS +#endif + +// The goal of this macro is to avoid evaluation of the arguments, but +// still have the compiler warn on problems inside... +#if !defined(CATCH_INTERNAL_IGNORE_BUT_WARN) +# define CATCH_INTERNAL_IGNORE_BUT_WARN(...) +#endif + +#if defined(__APPLE__) && defined(__apple_build_version__) && (__clang_major__ < 10) +# undef CATCH_INTERNAL_SUPPRESS_UNUSED_TEMPLATE_WARNINGS +#elif defined(__clang__) && (__clang_major__ < 5) +# undef CATCH_INTERNAL_SUPPRESS_UNUSED_TEMPLATE_WARNINGS +#endif + +#if !defined(CATCH_INTERNAL_SUPPRESS_UNUSED_TEMPLATE_WARNINGS) +# define CATCH_INTERNAL_SUPPRESS_UNUSED_TEMPLATE_WARNINGS +#endif + +#if defined(CATCH_CONFIG_DISABLE_EXCEPTIONS) +#define CATCH_TRY if ((true)) +#define CATCH_CATCH_ALL if ((false)) +#define CATCH_CATCH_ANON(type) if ((false)) +#else +#define CATCH_TRY try +#define CATCH_CATCH_ALL catch (...) +#define CATCH_CATCH_ANON(type) catch (type) +#endif + +#if defined(CATCH_INTERNAL_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR) && !defined(CATCH_CONFIG_NO_TRADITIONAL_MSVC_PREPROCESSOR) && !defined(CATCH_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR) +#define CATCH_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR +#endif + +// end catch_compiler_capabilities.h +#define INTERNAL_CATCH_UNIQUE_NAME_LINE2( name, line ) name##line +#define INTERNAL_CATCH_UNIQUE_NAME_LINE( name, line ) INTERNAL_CATCH_UNIQUE_NAME_LINE2( name, line ) +#ifdef CATCH_CONFIG_COUNTER +# define INTERNAL_CATCH_UNIQUE_NAME( name ) INTERNAL_CATCH_UNIQUE_NAME_LINE( name, __COUNTER__ ) +#else +# define INTERNAL_CATCH_UNIQUE_NAME( name ) INTERNAL_CATCH_UNIQUE_NAME_LINE( name, __LINE__ ) +#endif + +#include +#include +#include + +// We need a dummy global operator<< so we can bring it into Catch namespace later +struct Catch_global_namespace_dummy {}; +std::ostream& operator<<(std::ostream&, Catch_global_namespace_dummy); + +namespace Catch { + + struct CaseSensitive { enum Choice { + Yes, + No + }; }; + + class NonCopyable { + NonCopyable( NonCopyable const& ) = delete; + NonCopyable( NonCopyable && ) = delete; + NonCopyable& operator = ( NonCopyable const& ) = delete; + NonCopyable& operator = ( NonCopyable && ) = delete; + + protected: + NonCopyable(); + virtual ~NonCopyable(); + }; + + struct SourceLineInfo { + + SourceLineInfo() = delete; + SourceLineInfo( char const* _file, std::size_t _line ) noexcept + : file( _file ), + line( _line ) + {} + + SourceLineInfo( SourceLineInfo const& other ) = default; + SourceLineInfo& operator = ( SourceLineInfo const& ) = default; + SourceLineInfo( SourceLineInfo&& ) noexcept = default; + SourceLineInfo& operator = ( SourceLineInfo&& ) noexcept = default; + + bool empty() const noexcept { return file[0] == '\0'; } + bool operator == ( SourceLineInfo const& other ) const noexcept; + bool operator < ( SourceLineInfo const& other ) const noexcept; + + char const* file; + std::size_t line; + }; + + std::ostream& operator << ( std::ostream& os, SourceLineInfo const& info ); + + // Bring in operator<< from global namespace into Catch namespace + // This is necessary because the overload of operator<< above makes + // lookup stop at namespace Catch + using ::operator<<; + + // Use this in variadic streaming macros to allow + // >> +StreamEndStop + // as well as + // >> stuff +StreamEndStop + struct StreamEndStop { + std::string operator+() const; + }; + template + T const& operator + ( T const& value, StreamEndStop ) { + return value; + } +} + +#define CATCH_INTERNAL_LINEINFO \ + ::Catch::SourceLineInfo( __FILE__, static_cast( __LINE__ ) ) + +// end catch_common.h +namespace Catch { + + struct RegistrarForTagAliases { + RegistrarForTagAliases( char const* alias, char const* tag, SourceLineInfo const& lineInfo ); + }; + +} // end namespace Catch + +#define CATCH_REGISTER_TAG_ALIAS( alias, spec ) \ + CATCH_INTERNAL_START_WARNINGS_SUPPRESSION \ + CATCH_INTERNAL_SUPPRESS_GLOBALS_WARNINGS \ + namespace{ Catch::RegistrarForTagAliases INTERNAL_CATCH_UNIQUE_NAME( AutoRegisterTagAlias )( alias, spec, CATCH_INTERNAL_LINEINFO ); } \ + CATCH_INTERNAL_STOP_WARNINGS_SUPPRESSION + +// end catch_tag_alias_autoregistrar.h +// start catch_test_registry.h + +// start catch_interfaces_testcase.h + +#include + +namespace Catch { + + class TestSpec; + + struct ITestInvoker { + virtual void invoke () const = 0; + virtual ~ITestInvoker(); + }; + + class TestCase; + struct IConfig; + + struct ITestCaseRegistry { + virtual ~ITestCaseRegistry(); + virtual std::vector const& getAllTests() const = 0; + virtual std::vector const& getAllTestsSorted( IConfig const& config ) const = 0; + }; + + bool isThrowSafe( TestCase const& testCase, IConfig const& config ); + bool matchTest( TestCase const& testCase, TestSpec const& testSpec, IConfig const& config ); + std::vector filterTests( std::vector const& testCases, TestSpec const& testSpec, IConfig const& config ); + std::vector const& getAllTestCasesSorted( IConfig const& config ); + +} + +// end catch_interfaces_testcase.h +// start catch_stringref.h + +#include +#include +#include +#include + +namespace Catch { + + /// A non-owning string class (similar to the forthcoming std::string_view) + /// Note that, because a StringRef may be a substring of another string, + /// it may not be null terminated. + class StringRef { + public: + using size_type = std::size_t; + using const_iterator = const char*; + + private: + static constexpr char const* const s_empty = ""; + + char const* m_start = s_empty; + size_type m_size = 0; + + public: // construction + constexpr StringRef() noexcept = default; + + StringRef( char const* rawChars ) noexcept; + + constexpr StringRef( char const* rawChars, size_type size ) noexcept + : m_start( rawChars ), + m_size( size ) + {} + + StringRef( std::string const& stdString ) noexcept + : m_start( stdString.c_str() ), + m_size( stdString.size() ) + {} + + explicit operator std::string() const { + return std::string(m_start, m_size); + } + + public: // operators + auto operator == ( StringRef const& other ) const noexcept -> bool; + auto operator != (StringRef const& other) const noexcept -> bool { + return !(*this == other); + } + + auto operator[] ( size_type index ) const noexcept -> char { + assert(index < m_size); + return m_start[index]; + } + + public: // named queries + constexpr auto empty() const noexcept -> bool { + return m_size == 0; + } + constexpr auto size() const noexcept -> size_type { + return m_size; + } + + // Returns the current start pointer. If the StringRef is not + // null-terminated, throws std::domain_exception + auto c_str() const -> char const*; + + public: // substrings and searches + // Returns a substring of [start, start + length). + // If start + length > size(), then the substring is [start, size()). + // If start > size(), then the substring is empty. + auto substr( size_type start, size_type length ) const noexcept -> StringRef; + + // Returns the current start pointer. May not be null-terminated. + auto data() const noexcept -> char const*; + + constexpr auto isNullTerminated() const noexcept -> bool { + return m_start[m_size] == '\0'; + } + + public: // iterators + constexpr const_iterator begin() const { return m_start; } + constexpr const_iterator end() const { return m_start + m_size; } + }; + + auto operator += ( std::string& lhs, StringRef const& sr ) -> std::string&; + auto operator << ( std::ostream& os, StringRef const& sr ) -> std::ostream&; + + constexpr auto operator "" _sr( char const* rawChars, std::size_t size ) noexcept -> StringRef { + return StringRef( rawChars, size ); + } +} // namespace Catch + +constexpr auto operator "" _catch_sr( char const* rawChars, std::size_t size ) noexcept -> Catch::StringRef { + return Catch::StringRef( rawChars, size ); +} + +// end catch_stringref.h +// start catch_preprocessor.hpp + + +#define CATCH_RECURSION_LEVEL0(...) __VA_ARGS__ +#define CATCH_RECURSION_LEVEL1(...) CATCH_RECURSION_LEVEL0(CATCH_RECURSION_LEVEL0(CATCH_RECURSION_LEVEL0(__VA_ARGS__))) +#define CATCH_RECURSION_LEVEL2(...) CATCH_RECURSION_LEVEL1(CATCH_RECURSION_LEVEL1(CATCH_RECURSION_LEVEL1(__VA_ARGS__))) +#define CATCH_RECURSION_LEVEL3(...) CATCH_RECURSION_LEVEL2(CATCH_RECURSION_LEVEL2(CATCH_RECURSION_LEVEL2(__VA_ARGS__))) +#define CATCH_RECURSION_LEVEL4(...) CATCH_RECURSION_LEVEL3(CATCH_RECURSION_LEVEL3(CATCH_RECURSION_LEVEL3(__VA_ARGS__))) +#define CATCH_RECURSION_LEVEL5(...) CATCH_RECURSION_LEVEL4(CATCH_RECURSION_LEVEL4(CATCH_RECURSION_LEVEL4(__VA_ARGS__))) + +#ifdef CATCH_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR +#define INTERNAL_CATCH_EXPAND_VARGS(...) __VA_ARGS__ +// MSVC needs more evaluations +#define CATCH_RECURSION_LEVEL6(...) CATCH_RECURSION_LEVEL5(CATCH_RECURSION_LEVEL5(CATCH_RECURSION_LEVEL5(__VA_ARGS__))) +#define CATCH_RECURSE(...) CATCH_RECURSION_LEVEL6(CATCH_RECURSION_LEVEL6(__VA_ARGS__)) +#else +#define CATCH_RECURSE(...) CATCH_RECURSION_LEVEL5(__VA_ARGS__) +#endif + +#define CATCH_REC_END(...) +#define CATCH_REC_OUT + +#define CATCH_EMPTY() +#define CATCH_DEFER(id) id CATCH_EMPTY() + +#define CATCH_REC_GET_END2() 0, CATCH_REC_END +#define CATCH_REC_GET_END1(...) CATCH_REC_GET_END2 +#define CATCH_REC_GET_END(...) CATCH_REC_GET_END1 +#define CATCH_REC_NEXT0(test, next, ...) next CATCH_REC_OUT +#define CATCH_REC_NEXT1(test, next) CATCH_DEFER ( CATCH_REC_NEXT0 ) ( test, next, 0) +#define CATCH_REC_NEXT(test, next) CATCH_REC_NEXT1(CATCH_REC_GET_END test, next) + +#define CATCH_REC_LIST0(f, x, peek, ...) , f(x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST1) ) ( f, peek, __VA_ARGS__ ) +#define CATCH_REC_LIST1(f, x, peek, ...) , f(x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST0) ) ( f, peek, __VA_ARGS__ ) +#define CATCH_REC_LIST2(f, x, peek, ...) f(x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST1) ) ( f, peek, __VA_ARGS__ ) + +#define CATCH_REC_LIST0_UD(f, userdata, x, peek, ...) , f(userdata, x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST1_UD) ) ( f, userdata, peek, __VA_ARGS__ ) +#define CATCH_REC_LIST1_UD(f, userdata, x, peek, ...) , f(userdata, x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST0_UD) ) ( f, userdata, peek, __VA_ARGS__ ) +#define CATCH_REC_LIST2_UD(f, userdata, x, peek, ...) f(userdata, x) CATCH_DEFER ( CATCH_REC_NEXT(peek, CATCH_REC_LIST1_UD) ) ( f, userdata, peek, __VA_ARGS__ ) + +// Applies the function macro `f` to each of the remaining parameters, inserts commas between the results, +// and passes userdata as the first parameter to each invocation, +// e.g. CATCH_REC_LIST_UD(f, x, a, b, c) evaluates to f(x, a), f(x, b), f(x, c) +#define CATCH_REC_LIST_UD(f, userdata, ...) CATCH_RECURSE(CATCH_REC_LIST2_UD(f, userdata, __VA_ARGS__, ()()(), ()()(), ()()(), 0)) + +#define CATCH_REC_LIST(f, ...) CATCH_RECURSE(CATCH_REC_LIST2(f, __VA_ARGS__, ()()(), ()()(), ()()(), 0)) + +#define INTERNAL_CATCH_EXPAND1(param) INTERNAL_CATCH_EXPAND2(param) +#define INTERNAL_CATCH_EXPAND2(...) INTERNAL_CATCH_NO## __VA_ARGS__ +#define INTERNAL_CATCH_DEF(...) INTERNAL_CATCH_DEF __VA_ARGS__ +#define INTERNAL_CATCH_NOINTERNAL_CATCH_DEF +#define INTERNAL_CATCH_STRINGIZE(...) INTERNAL_CATCH_STRINGIZE2(__VA_ARGS__) +#ifndef CATCH_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR +#define INTERNAL_CATCH_STRINGIZE2(...) #__VA_ARGS__ +#define INTERNAL_CATCH_STRINGIZE_WITHOUT_PARENS(param) INTERNAL_CATCH_STRINGIZE(INTERNAL_CATCH_REMOVE_PARENS(param)) +#else +// MSVC is adding extra space and needs another indirection to expand INTERNAL_CATCH_NOINTERNAL_CATCH_DEF +#define INTERNAL_CATCH_STRINGIZE2(...) INTERNAL_CATCH_STRINGIZE3(__VA_ARGS__) +#define INTERNAL_CATCH_STRINGIZE3(...) #__VA_ARGS__ +#define INTERNAL_CATCH_STRINGIZE_WITHOUT_PARENS(param) (INTERNAL_CATCH_STRINGIZE(INTERNAL_CATCH_REMOVE_PARENS(param)) + 1) +#endif + +#define INTERNAL_CATCH_MAKE_NAMESPACE2(...) ns_##__VA_ARGS__ +#define INTERNAL_CATCH_MAKE_NAMESPACE(name) INTERNAL_CATCH_MAKE_NAMESPACE2(name) + +#define INTERNAL_CATCH_REMOVE_PARENS(...) INTERNAL_CATCH_EXPAND1(INTERNAL_CATCH_DEF __VA_ARGS__) + +#ifndef CATCH_CONFIG_TRADITIONAL_MSVC_PREPROCESSOR +#define INTERNAL_CATCH_MAKE_TYPE_LIST2(...) decltype(get_wrapper()) +#define INTERNAL_CATCH_MAKE_TYPE_LIST(...) INTERNAL_CATCH_MAKE_TYPE_LIST2(INTERNAL_CATCH_REMOVE_PARENS(__VA_ARGS__)) +#else +#define INTERNAL_CATCH_MAKE_TYPE_LIST2(...) INTERNAL_CATCH_EXPAND_VARGS(decltype(get_wrapper())) +#define INTERNAL_CATCH_MAKE_TYPE_LIST(...) INTERNAL_CATCH_EXPAND_VARGS(INTERNAL_CATCH_MAKE_TYPE_LIST2(INTERNAL_CATCH_REMOVE_PARENS(__VA_ARGS__))) +#endif + +#define INTERNAL_CATCH_MAKE_TYPE_LISTS_FROM_TYPES(...)\ + CATCH_REC_LIST(INTERNAL_CATCH_MAKE_TYPE_LIST,__VA_ARGS__) + +#define INTERNAL_CATCH_REMOVE_PARENS_1_ARG(_0) INTERNAL_CATCH_REMOVE_PARENS(_0) +#define INTERNAL_CATCH_REMOVE_PARENS_2_ARG(_0, _1) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_1_ARG(_1) +#define INTERNAL_CATCH_REMOVE_PARENS_3_ARG(_0, _1, _2) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_2_ARG(_1, _2) +#define INTERNAL_CATCH_REMOVE_PARENS_4_ARG(_0, _1, _2, _3) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_3_ARG(_1, _2, _3) +#define INTERNAL_CATCH_REMOVE_PARENS_5_ARG(_0, _1, _2, _3, _4) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_4_ARG(_1, _2, _3, _4) +#define INTERNAL_CATCH_REMOVE_PARENS_6_ARG(_0, _1, _2, _3, _4, _5) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_5_ARG(_1, _2, _3, _4, _5) +#define INTERNAL_CATCH_REMOVE_PARENS_7_ARG(_0, _1, _2, _3, _4, _5, _6) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_6_ARG(_1, _2, _3, _4, _5, _6) +#define INTERNAL_CATCH_REMOVE_PARENS_8_ARG(_0, _1, _2, _3, _4, _5, _6, _7) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_7_ARG(_1, _2, _3, _4, _5, _6, _7) +#define INTERNAL_CATCH_REMOVE_PARENS_9_ARG(_0, _1, _2, _3, _4, _5, _6, _7, _8) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_8_ARG(_1, _2, _3, _4, _5, _6, _7, _8) +#define INTERNAL_CATCH_REMOVE_PARENS_10_ARG(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_9_ARG(_1, _2, _3, _4, _5, _6, _7, _8, _9) +#define INTERNAL_CATCH_REMOVE_PARENS_11_ARG(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10) INTERNAL_CATCH_REMOVE_PARENS(_0), INTERNAL_CATCH_REMOVE_PARENS_10_ARG(_1, _2, _3, _4, _5, _6, _7, _8, _9, _10) + +#define INTERNAL_CATCH_VA_NARGS_IMPL(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, N, ...) N + +#define INTERNAL_CATCH_TYPE_GEN\ + template struct TypeList {};\ + template\ + constexpr auto get_wrapper() noexcept -> TypeList { return {}; }\ + template class...> struct TemplateTypeList{};\ + template class...Cs>\ + constexpr auto get_wrapper() noexcept -> TemplateTypeList { return {}; }\ + template\ + struct append;\ + template\ + struct rewrap;\ + template class, typename...>\ + struct create;\ + template class, typename>\ + struct convert;\ + \ + template \ + struct append { using type = T; };\ + template< template class L1, typename...E1, template class L2, typename...E2, typename...Rest>\ + struct append, L2, Rest...> { using type = typename append, Rest...>::type; };\ + template< template class L1, typename...E1, typename...Rest>\ + struct append, TypeList, Rest...> { using type = L1; };\ + \ + template< template class Container, template class List, typename...elems>\ + struct rewrap, List> { using type = TypeList>; };\ + template< template class Container, template class List, class...Elems, typename...Elements>\ + struct rewrap, List, Elements...> { using type = typename append>, typename rewrap, Elements...>::type>::type; };\ + \ + template