Skip to content

TESTS.md

Latest commit

 

History

History
402 lines (342 loc) · 144 KB

TESTS.md

File metadata and controls

402 lines (342 loc) · 144 KB
Raw

algorithm_benchmark

Algorithm Benchmark is a set of benchmarks aimed at measuring the performance of realistic worloads.

Test name Description Params L0 OCL
Heat3D A 3D heat-equation solving benchmark that overlaps IPC data transfers (nearest-neighbor halo exchange) and GPU compute kernels.Measures multi-process concurrent kernel execution and IPC memory transfer performance on a single device. Linux-only.
  • --meshLength Number of mesh points along each of the X-Y-Z directions
  • --subDomainX Number of sub-domains in the X-direction
  • --subDomainY Number of sub-domains in the Y-direction
  • --subDomainZ Number of sub-domains in the Z-direction
  • --timesteps Number of simulation timesteps
 ✔️ ❌

api_overhead_benchmark

Api Overhead Benchmark is a set of tests aimed at measuring CPU-side execution duration of compute API calls.

Test name Description Params L0 OCL
AppendImageCopyFromMemoryExt measures time spent in zeCommandListAppendImageCopyFromMemoryExt on CPU. The destination is an RGBA float image and the source row and slice pitches are 0, so the driver computes them.
  • --appendCount Number of appends to run
  • --event Pass output event to the append call (0 or 1)
  • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
  • --size Size of the destination image
  • --src Placement of source memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
  • --wholeImage Copy the whole image by passing no destination region, as opposed to a region of a single pixel (0 or 1)
 ✔️ ❌
AppendImageCopyToMemoryExt measures time spent in zeCommandListAppendImageCopyToMemoryExt on CPU. The source is an RGBA float image and the destination row and slice pitches are 0, so the driver computes them.
  • --appendCount Number of appends to run
  • --dst Placement of destination memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
  • --event Pass output event to the append call (0 or 1)
  • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
  • --size Size of the source image
  • --wholeImage Copy the whole image by passing no source region, as opposed to a region of a single pixel (0 or 1)
 ✔️ ❌
AppendKernelWithMappedTimestampEvent Measures CPU time of zeCommandListAppendLaunchKernel with timestamp event.
  • --useMappedTimestamp Use ZE_EVENT_POOL_FLAG_KERNEL_MAPPED_TIMESTAMP instead of ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP. (0 or 1)
 ✔️ ❌
AppendLaunchCooperativeKernel measures time spent in zeCommandListAppendLaunchCooperativeKernel on CPU.
  • --appendCount Number of appends to run
  • --event Pass output event to the enqueue call (0 or 1)
  • --wgc Workgroup count
  • --wgs Workgroup size, pass 0 to make the driver calculate it during enqueue
 ✔️ ❌
AppendLaunchKernel measures time spent in zeCommandListAppendLaunchKernel on CPU.
  • --appendCount Number of appends to run
  • --event Pass output event to the enqueue call (0 or 1)
  • --wgc Workgroup count
  • --wgs Workgroup size, pass 0 to make the driver calculate it during enqueue
 ✔️ ❌
AppendLaunchKernelWithParameters measures time spent in zeCommandListAppendLaunchKernelWithParameters on CPU.
  • --appendCount Number of appends to run
  • --cooperative Pass cooperative descriptor via pNext (0 or 1)
  • --event Pass output event to the enqueue call (0 or 1)
  • --wgc Workgroup count
  • --wgs Workgroup size, pass 0 to make the driver calculate it during enqueue
 ✔️ ❌
AppendMarkerExp measures time spent in zetCommandListAppendMarkerExp on CPU. The first metric group of the device that can generate markers is activated. Requires ZET_ENABLE_METRICS=1.
  • --appendCount Number of appends to run
 ✔️ ❌
AppendMemCopy measures time spent in zeCommandListAppendMemoryCopy on CPU.
  • --appendCount Number of appends to run
  • --dst Placement of destination memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
  • --event Pass output event to the append call (0 or 1)
  • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
  • --size Size of memory to copy
  • --src Placement of source memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
 ✔️ ❌
AppendMetricMemoryBarrier measures time spent in zetCommandListAppendMetricMemoryBarrier on CPU. The barriers are appended between the begin and the end of a metric query from the first event based metric group of the device. Requires ZET_ENABLE_METRICS=1 and access to the device performance counters.
  • --appendCount Number of appends to run
 ✔️ ❌
AppendMetricQueryBegin measures time spent in zetCommandListAppendMetricQueryBegin on CPU. The metric query pool uses the first event based metric group of the device. Requires ZET_ENABLE_METRICS=1 and access to the device performance counters.
  • --appendCount Number of appends to run, each with its own metric query
 ✔️ ❌
AppendMetricQueryEnd measures time spent in zetCommandListAppendMetricQueryEnd on CPU. The metric query pool uses the first event based metric group of the device. Requires ZET_ENABLE_METRICS=1 and access to the device performance counters.
  • --appendCount Number of appends to run, each with its own metric query
  • --event Pass a signal event from an event pool to the append call (0 or 1)
 ✔️ ❌
AppendMetricStreamerMarker measures time spent in zetCommandListAppendMetricStreamerMarker on CPU. The metric streamer is opened on the first time based metric group of the device that accepts a streamer marker. Requires ZET_ENABLE_METRICS=1 and access to the device performance counters.
  • --appendCount Number of appends to run
 ✔️ ❌
AppendSignalEvent measures time spent in zeCommandListAppendSignalEvent on CPU.
  • --appendCount Number of appends to run
  • --counterBasedEvents Use counter-based events instead of events from an event pool (0 or 1)
  • --timestampEvents Create events with kernel timestamp support (0 or 1)
 ✔️ ❌
AppendWaitOnEventsImmediate Measures time spent to zeCommandListAppendWaitOnEvents using immediate command list.
  • --eventSignaled Event is already signaled before zeCommandListAppendWaitOnEvents call (0 or 1)
  • --ioq Use In order queue (0 or 1)
 ✔️ ❌
CommandListGetContextHandle measures time spent in zeCommandListGetContextHandle call on CPU.
  • --useImmediate Query an immediate command list instead of a regular one (0 or 1)
 ✔️ ❌
CommandListGetDeviceHandle measures time spent in zeCommandListGetDeviceHandle call on CPU.
  • --useImmediate Query an immediate command list instead of a regular one (0 or 1)
 ✔️ ❌
CommandListGetNextCommandId measures time spent in zeCommandListGetNextCommandIdExp call on CPU. Supported mutableCommandFlag values: All, KernelArguments, GroupCount, GroupSize, GlobalOffset, SignalEvent, WaitEvents, GraphArguments.
  • --mutableCommandFlag Specifies the mutable command flag to be used in the test. Supported values: All, KernelArguments, GroupCount, GroupSize, GlobalOffset, SignalEvent, WaitEvents, GraphArguments. (All or KernelArguments or GroupCount or GroupSize or GlobalOffset or SignalEvent or WaitEvents or KernelInstruction or GraphArguments)
 ✔️ ❌
CommandListGetNextCommandIdWithKernels measures time spent in zeCommandListGetNextCommandIdWithKernelsExp call on CPU. Supported mutableCommandFlag values: All, KernelArguments, GroupCount, GroupSize, GlobalOffset, SignalEvent, WaitEvents, KernelInstruction, GraphArguments.
  • --mutableCommandFlag Specifies the mutable command flag to be used in the test. Supported values: All, KernelArguments, GroupCount, GroupSize, GlobalOffset, SignalEvent, WaitEvents, KernelInstruction, GraphArguments. (All or KernelArguments or GroupCount or GroupSize or GlobalOffset or SignalEvent or WaitEvents or KernelInstruction or GraphArguments)
  • --numKernels Number of kernels that user can switch between using zeCommandListUpdateMutableCommandKernelsExp call.
 ✔️ ❌
CommandListGetOrdinal measures time spent in zeCommandListGetOrdinal call on CPU.
  • --useImmediate Query an immediate command list instead of a regular one (0 or 1)
 ✔️ ❌
CommandListHostSynchronize measures CPU time spent in zeCommandListHostSynchronize. Optionally, adds an event-signalling barrierand waits for the event, before calling zeCommandListHostSynchronize
  • --UseBarrierBeforeSync Append an event-signalling-barrier before synchronization (0 or 1)
 ✔️ ❌
CommandListImmediateGetIndex measures time spent in zeCommandListImmediateGetIndex call on CPU.
    		✔️ ❌
    CommandListIsImmediate measures time spent in zeCommandListIsImmediate call on CPU.
    • --useImmediate Query an immediate command list instead of a regular one (0 or 1)
    		 ✔️ ❌
    CommandListUpdateMutableCommandKernels measures time spent in zeCommandListUpdateMutableCommandKernelsExp call on CPU. A single mutable kernel command is appended with one initial kernel selected from a list of N candidates. Each iteration swaps the active kernel to the next one in the list.
    • --numKernels Number of kernels registered with the mutable command that can be swapped in via zeCommandListUpdateMutableCommandKernelsExp.
    		 ✔️ ❌
    CommandListUpdateMutableCommandSignalEvent measures time spent in zeCommandListUpdateMutableCommandSignalEventExp call on CPU. A single mutable kernel command is appended with an initial signal event. Each iteration swaps the command's signal event to an alternate handle.
      		✔️ ❌
      CommandListUpdateMutableCommandWaitEvents measures time spent in zeCommandListUpdateMutableCommandWaitEventsExp call on CPU. A single mutable kernel command is appended with N wait events. Each iteration replaces the wait event list with an alternate set of N events.
      • --numWaitEvents Number of wait events the mutable kernel command was appended with and that are swapped each iteration.
      		 ✔️ ❌
      CommandListUpdateMutableCommands measures time spent in zeCommandListUpdateMutableCommandsExp call on CPU. Mutates a single previously-appended kernel command via mutation descriptors chained on pNext. Supported mutableCommandFlag values: KernelArguments, GroupCount, GroupSize, GlobalOffset, All.
      • --mutableCommandFlag Specifies which mutation type(s) to apply via zeCommandListUpdateMutableCommandsExp. Supported values: KernelArguments, GroupCount, GroupSize, GlobalOffset, All. 'All' chains kernelArgs + groupCount + groupSize + globalOffset for the single mutable command. (All or KernelArguments or GroupCount or GroupSize or GlobalOffset or SignalEvent or WaitEvents or KernelInstruction or GraphArguments)
      		 ✔️ ❌
      CreateBuffer measures time spent in clCreateBuffer on CPU.
      • --allocateAll Free buffers at the end of test, as opposed to freeing between iterations. Should disallow resource reuse (0 or 1)
      • --bufferSize Buffer size
      • --copyHostPtr CL_MEM_COPY_HOST_PTR flag (0 or 1)
      • --forceHostMemoryIntel CL_MEM_FORCE_HOST_MEMORY_INTEL flag (0 or 1)
      • --readOnly Read only buffer (0 or 1)
      		 ❌ ✔️
      CreateCommandList measures time spent in zeCommandListCreate on CPU.
      • --CmdListCount Number of cmdlists to create
      • --CopyOnly Create copy only cmdlist (0 or 1)
      		 ✔️ ❌
      CreateCommandListImmediate measures time spent in zeCommandListCreateImmediate on CPU.
      • --CmdListCount Number of cmdlists to create
      • --ioq Use In order queue (0 or 1)
      		 ✔️ ❌
      DestroyCommandList measures time spent in zeCommandListDestroy on CPU.
      • --CmdListCount Number of cmdlists to destroy
      		 ✔️ ❌
      DestroyCommandListImmediate measures time spent in zeCommandListDestroy on CPU, for immediate cmdlist.
      • --CmdListCount Number of immediate cmdlists to create
      • --ioq Use In order queue (0 or 1)
      		 ✔️ ❌
      DriverGet measures time spent in driver get call on CPU.
      • --getDriverCount Whether to measure driver count or driver get (0 or 1)
      		 ✔️ ❌
      DriverGetApiVersion measures time spent in zeDriverGetApiVersion call on CPU.
        		✔️ ❌
        DriverGetProperties measures time spent in zeDriverGetProperties call on CPU.
          		✔️ ❌
          EnqueueNdrNullLws measures time spent in clEnqueueNDRangeKernel on CPU. Null LWS is provided, which causes driver to calculate it
          • --event Pass output event to the enqueue call (0 or 1)
          • --gws Global work size
          • --ooq Use out of order queue (0 or 1)
          • --profiling Creating a profiling queue (0 or 1)
          		 ❌ ✔️
          EnqueueNdrTime measures time spent in clEnqueueNDRangeKernel on CPU.
          • --event Pass output event to the enqueue call (0 or 1)
          • --ooq Use out of order queue (0 or 1)
          • --profiling Creating a profiling queue (0 or 1)
          • --wgc Workgroup count
          • --wgs Workgroup size
          		 ❌ ✔️
          EventCreation measures time spent to create event
          • --eventCount Number of events to create
          • --hostVisible Event will set host visible flag (0 or 1)
          • --signal Type of signal scope (subdevice or device or host or none)
          • --useCounterBasedEvents Event will be of counter based type (0 or 1)
          • --useProfiling Event will use profiling (0 or 1)
          • --wait Type of wait scope (subdevice or device or host or none)
          		 ✔️ ❌
          EventDestroy measures time spent to destroy completed event
          • --eventCount Number of events to create
          • --hostVisible Event will set host visible flag (0 or 1)
          • --signal Type of signal scope (subdevice or device or host or none)
          • --useCounterBasedEvents Event will be of counter based type (0 or 1)
          • --useProfiling Event will use profiling (0 or 1)
          • --wait Type of wait scope (subdevice or device or host or none)
          		 ✔️ ❌
          EventQueryKernelTimestamp Measures time spent in zeEventQueryKernelTimestamp for a completed counter based event with device timestamps. First is the cost of the first query, Repeated is the per call cost of the queries that follow it.
          • --queryCount Number of queries performed after the first one
          • --useImmediate Signal the event from an immediate command list instead of a regular one (0 or 1)
          		 ✔️ ❌
          EventQueryProfilingData measures time spent querying profiling data from events
          • --eventCount Number of events to query profiling data from
          • --splitCount Number of internal timestamp splits per operation (1, 2, or 3)
          		 ❌ ✔️
          EventQueryStatus Measures time spent to query event status
          • --eventSignaled Event will be set as signaled (0 or 1)
          		 ✔️ ❌
          EventReuseImmCmdlistInProgress measures CPU-side time of zeCommandListAppendLaunchKernel when reusing a signal event on an immediate command list that is currently in progress.
          • --kernelExecutionTime Approximately how long the stall kernel executes, in us
          		 ✔️ ❌
          ExecImmCopy measures time spent in appending memory copy for immediate command list on CPU with Copy Queue.
          • --CopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
          • --IsCopyOnly If true, Copy Engine is selected. If false, Compute Engine is selected (0 or 1)
          • --MeasureCompletionTime Measures time taken to complete the submission (default is to measure only Immediate call) (0 or 1)
          • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
          • --ioq Use In order queue (0 or 1)
          • --size Size of the buffer
          • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
          		 ✔️ ❌
          ExecImmediate measures time spent in appending launch kernel for immediate command list on CPU.
          • --BarrierSynchro Uses barrier synchronization instead of waiting for event from last kernel (0 or 1)
          • --CallsCount amount of calls that is being meassured
          • --EventSync If true, use events to synchronize with host. If false, use zeCommandListHostSynchronize (0 or 1)
          • --KernelExecTime How long a single kernel executes, in us
          • --MeasureCompletion Measures time taken to complete the submission (default is to measure only Immediate call) (0 or 1)
          • --Profiling Pass a profiling ze_event_t to the API call (0 or 1)
          • --ioq Use In order queue (0 or 1)
          		 ✔️ ❌
          ExecImmediateMultiKernel measures time spent in executing multiple instances of two different kernels with immediate command list on CPU
          • --AfterBarrierCnt Adds certain number of kernels after Barrier, Default is 2
          • --Barrier Add a Barrier after certain number of Kernel launches, number of kernels before barrier is controlled by numKernelsBeforeBarrier (0 or 1)
          • --BeforeBarrierCnt Adds certain number of kernels prior to Barrier, Default is 2
          • --CallsCount amount of calls that is being measured
          • --ExecTime Approximately how long a single kernel executes, in us
          • --ioq Use In order queue (0 or 1)
          		 ✔️ ❌
          ExecuteCommandList measures time spent in zeCommandQueueExecuteCommandLists on CPU.
          • --UseFence Pass a non-null ze_fence_handle_t to the API call (0 or 1)
          • --measureCompletionTime Measures time taken to complete the submission (default is to measure only Execute call) (0 or 1)
          		 ✔️ ❌
          ExecuteCommandListForCopyEngine measures CPU time spent in zeCommandQueueExecuteCommandLists for copy-only path
          • --UseFence Pass a non-null ze_fence_handle_t to the API call (0 or 1)
          • --measureCompletionTime Measures time taken to complete the submission (default is to measure only Execute call) (0 or 1)
          		 ✔️ ❌
          ExecuteCommandListWithFenceCreate measures time spent in zeFenceCreate on CPU when fences are used.
            		✔️ ❌
            ExecuteCommandListWithFenceDestroy measures time spent in zeFenceDestroy on CPU when fences are used.
              		✔️ ❌
              ExecuteCommandListWithFenceUsage measures time spent in zeCommandQueueExecuteCommandLists and zeFenceSynchronize on CPU when fences are used.
                		✔️ ❌
                ExecuteCommandListWithIndirectAccess measures time spent in zeCommandQueueExecuteCommandLists on CPU when indirect allocations are accessed.
                • --AllocateMemory If set then prior to measurement new allocation is done and made resident. (0 or 1)
                • --AmountOfIndirectAllocations Amount of indirect allocations that are present in system
                		 ✔️ ❌
                ExecuteCommandListWithIndirectArguments measures time spent in zeCommandQueueExecuteCommandLists on CPU when indirect allocations are used.
                • --AmountOfIndirectAllocations Amount of indirect allocations that are present in system
                • --placement Placement of the indirect allocations (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                		 ✔️ ❌
                ExecuteRegularCommandListWithImmediate measures time spent in zeCommandListImmediateAppendCommandListsExp on CPU.
                • --ctrBasedEvents use counter based events for in order (0 or 1)
                • --event Pass output event to the enqueue call (0 or 1)
                • --ioq use in order queue/command list (0 or 1)
                • --measureCompletionTime Measures time taken to complete the submission (0 or 1)
                • --useProfiling Event will use profiling (0 or 1)
                • --waitEvent use preceeding operation to wait on (0 or 1)
                		 ✔️ ❌
                FlushTime measures time spent in clEnqueueNDRangeKernel on CPU.
                • --event Pass output event to the enqueue call (0 or 1)
                • --flushCount Count of flushes to measure
                • --ooq Use out of order queue (0 or 1)
                • --wgc Workgroup count
                • --wgs Workgroup size, pass 0 to make the driver calculate it during enqueue
                		 ❌ ✔️
                GetMemoryProperties measures time spent in zeMemGetAllocProperties on CPU when driver is queried for memory properties.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                		 ✔️ ❌
                GetMemoryPropertiesWithModifiedAllocations measures time spent in zeMemGetAllocProperties on CPU, when allocations are modified between each iteration.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                		 ✔️ ❌
                GetMemoryPropertiesWithOffsetedPointer measures time spent in zeMemGetAllocProperties on CPU when the pointer passed is an offset from the base address.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                		 ✔️ ❌
                ImageRecreation measures time used to create a temporary image, write to it on the GPU, wait for the write and release the image, as in a frame loop that allocates a new image for every frame.
                • --size Size of the image
                • --writeWholeImage Write the whole image on the GPU, as opposed to a single pixel (0 or 1)
                		 ✔️ ✔️
                InOrderWaitAppend Measures time spent to append wait on event command to in-order command list.
                • --counterBasedEvents Use Counter Based Events (0 or 1)
                • --isCompleted Is the event completed before appending (0 or 1)
                • --kernelExecutionTime How long a single kernel executes, in us in non-completed mode
                		 ✔️ ❌
                InOrderWaitlistAppend Measures time spent to append a kernel to an in-order command list, when an event signaled by the very same command list is passed on the wait list of that append.
                • --counterBasedEvents Use Counter Based Events (0 or 1)
                • --isCompleted Is the event completed before appending (0 or 1)
                • --kernelExecutionTime How long the signaling kernel executes, in us in non-completed mode
                		 ✔️ ❌
                KernelSetArgumentValueImmediate measures time spent in zeKernelSetArgumentValue for immediate arguments on CPU.
                • --argSize Kernel argument size in bytes
                • --differentValues Use different values for arguments each iteration (0 or 1)
                		 ✔️ ❌
                LifecycleCommandList measures time spent on CPU to create, close, execute, synchronize and destroy CmdListCount command lists, one at a time.
                • --CmdListCount Number of cmdlists to create
                • --CopyOnly Create copy only cmdlist (0 or 1)
                		 ✔️ ❌
                MemGetIpcHandle measures time spent in zeMemGetIpcHandle on CPU.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                		 ✔️ ❌
                MemOpenIpcHandle measures time spent in zeMemOpenIpcHandle on CPU.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                		 ✔️ ❌
                MemPutIpcHandle measures time spent in zeMemPutIpcHandle on CPU.
                • --AmountOfUsmAllocations Amount of USM allocations that are present in system
                • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                		 ✔️ ❌
                ModuleCreateSpv measures time spent in zeModuleCreate for .spv kernel on CPU.
                • --kernelName Path to Kernel .spv file
                		 ✔️ ❌
                MultiArgumentKernelTime measures time spent in clEnqueueNDRangeKernel on CPU for kernels with multiple arguments.
                • --argumentCount argument count in a kernel, supported values 1,4,8,16,32,64
                • --count how many kernel launches are measured
                • --exec measure execute as well (0 or 1)
                • --groupCount total amount of work groups
                • --ids control whether kernel uses global ids (0 or 1)
                • --lws local work size
                • --measureSetArg control whether setKernelArgSvmPointer is measured or not (0 or 1)
                • --newApi utilize zeCommandListAppendLaunchKernelWithArguments for L0 submissions, on Vulkan writes all arguments of a kernel with a single vkUpdateDescriptorSets call instead of one call per argument (0 or 1)
                • --reverseOrder set kernel arguments in reverse order (0 or 1)
                		 ✔️ ✔️
                PhysicalMemCreate measures time spent in zePhysicalMemCreate on CPU.
                • --reserveSize Size in bytes to be reserved
                		 ✔️ ❌
                PhysicalMemDestroy measures time spent in zePhysicalMemDestroy on CPU.
                  		✔️ ❌
                  ResetCommandList measures time spent in zeCommandListReset on CPU.
                  • --CopyOnly Create copy only cmdlist (0 or 1)
                  • --size Size of the buffer
                  • --sourcePlacement Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                  		 ✔️ ❌
                  SetKernelArgSvmPointer measures time spent in clSetKernelArgSVMPointer on CPU.
                  • --allocationSize Size of svm allocations, in bytes
                  • --allocationsCount Number of allocations
                  • --reallocate Allocations will be freed and allocated again between setKernelArgs (0 or 1)
                  		 ✔️ ✔️
                  SetKernelGroupSize measures time spent in zeKernelSetGroupSize on CPU.
                  • --asymmetricLocalWorkSize Use asymmetric local workSize (0 or 1)
                  		 ✔️ ❌
                  SubmitKernel measures time spent in submitting a kernel to a SYCL (or SYCL-like) queue on CPU.
                  • --Ioq Create the queue with the in_order property (0 or 1)
                  • --KernelExecTime Approximately how long a single kernel executes, in us
                  • --MeasureCompletion Measures time taken to complete the submission (default is to measure only submit calls) (0 or 1)
                  • --NumKernels Number of kernels to submit to the queue
                  • --Profiling Create the queue with the enable_profiling property (0 or 1)
                  • --UseEvents Use events when enqueuing kernels. When false, SYCL will use eventless enqueue functions. (0 or 1)
                  		 ✔️ ✔️
                  UsmCopyReservedMemory measures time spent in zeCommandListAppendMemoryCopy on CPU on an immediate command list, when the source and destination are inside a virtual memory reservation that is mapped with the given number of zeVirtualMemMap calls.
                  • --chunkSize Size of each mapped chunk, rounded up to the virtual memory page size
                  • --mapCount Number of zeVirtualMemMap calls, each mapping its own physical memory chunk into the reservation
                  • --measureCompletionTime Measures time taken to complete the copy (default is to measure only the append call) (0 or 1)
                  • --size Size of memory to copy. Source and destination are placed back to back at the start of the reservation
                  		 ✔️ ❌
                  UsmMemoryAllocation measures time spent in USM memory allocation APIs using the specified strategy.
                  • --measureMode Specifies which APIs to measure (Allocate or Free or Both)
                  • --size Size to allocate
                  • --strategy Strategy used to allocate memory (Static is not applicable here, Async is supported only by UR and SYCL) (async or static or sync)
                  • --type Type of memory being allocated (Device or Host or Shared)
                  		 ✔️ ❌
                  VirtualMemFree measures time spent in zeVirtualMemFree on CPU.
                  • --freeSize Size in bytes to be freed
                  		 ✔️ ❌
                  VirtualMemGetAccessAttrib measures time spent in zeVirtualMemGetAccessAttribute on CPU.
                  • --size Size in bytes to get the access attribute
                  		 ✔️ ❌
                  VirtualMemMap measures time spent in zeVirtualMemMap on CPU.
                  • --accessType Access type. Either 'ReadWrite' or 'ReadOnly'
                  • --poolSize Upper bound on virtual/physical memory reservations held at once. Each iteration, warmup included, maps its own reservation, so iterations plus warmup iterations must not exceed it
                  • --reserveSize Size in bytes to be reserved
                  • --useOffset Use offset to map into physical memory (0 or 1)
                  		 ✔️ ❌
                  VirtualMemQueryPageSize measures time spent in zeVirtualMemQueryPageSize on CPU.
                    		✔️ ❌
                    VirtualMemReserve measures time spent in zeVirtualMemReserve on CPU.
                    • --reserveSize Size in bytes to be reserved
                    • --useNull Flag to decide whether Null to be used for start of region (0 or 1)
                    		 ✔️ ❌
                    VirtualMemSetAccessAttrib measures time spent in zeVirtualMemSetAccessAttribute on CPU.
                    • --accessType Access type to set. Either 'ReadWrite', 'ReadOnly' or 'None'
                    • --size Size in bytes to set the access attribute
                    		 ✔️ ❌
                    VirtualMemUnMap measures time spent in zeVirtualMemUnMap on CPU.
                    • --poolSize Number of virtual/physical memory reservations to pool and cycle through
                    • --reserveSize Size in bytes to be unmapped
                    		 ✔️ ❌

                    atomic_benchmark

                    Atomic Benchmark is a set of tests aimed at measuring performance of atomic operations inside kernels.

                    Test name Description Params L0 OCL
                    OneAtomic enqueues kernel performing an atomic operation on a single address
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️
                    OneAtomicExplicit enqueues kernel performing an atomic operation on a single address using OpenCL 2.0 Atomics with explicit memory order and scope
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --order Memory order of an atomic operation (relaxed or acquire or release or acq_rel or seq_cst)
                    • --scope Memory scope of an atomic operation (Workgroup or Device)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️
                    OneLocalAtomic enqueues kernel performing an atomic operation on a single location placed in SLM
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgs Work group size
                    		 ✔️ ✔️
                    OneLocalAtomicExplicit enqueues kernel performing an atomic operation on a single location placed in SLM using OpenCL 2.0 Atomics with explicit memory order and scope
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --order Memory order of an atomic operation (relaxed or acquire or release or acq_rel or seq_cst)
                    • --scope Memory scope of an atomic operation (Workgroup or Device)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgs Work group size
                    		 ✔️ ✔️
                    SeparateAtomics enqueues kernel performing an atomic operation on different addresses
                    • --atomicsPerCacheline Number of used addresses occupying a single cacheline (this causes operations to be serialized)
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️
                    SeparateAtomicsExplicit enqueues kernel performing an atomic operation on different addresses
                    • --atomicsPerCacheline Number of used addresses occupying a single cacheline (this causes operations to be serialized)
                    • --op Atomic operation to perform (Add or Sub or Xchg or CmpXchg or Inc or Dec or Min or Max or And or Or or Xor)
                    • --order Memory order of an atomic operation (relaxed or acquire or release or acq_rel or seq_cst)
                    • --scope Memory scope of an atomic operation (Workgroup or Device)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️

                    compute_benchmark

                    Compute Benchmark is a set of tests aimed at measuring floating point compute throughput of GPU kernels.

                    Test name Description Params L0 OCL
                    SinglePrecisionPerformance measures single precision floating point throughput in GFLOPS by executing FMA operations in a kernel, dispatch size is determined automatically from device capabilities
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️

                    cpu_efficiency_benchmark

                    CPU Efficiency Benchmark measures CPU cost, utilization, and latency tradeoffs of host-side wait and synchronization paths.

                    Test name Description Params L0 OCL
                    EventHostSynchronize Measures zeEventHostSynchronize latency, thread and process CPU time, and CPU utilization for work submitted to an immediate command list
                    • --batchSize Number of zeEventHostSynchronize calls measured per result
                    • --inOrderQueue Use an in-order immediate command list with counter-based events (0 or 1)
                    • --kernelExecutionTime Approximately how long a single kernel executes, in us
                    • --useKernelTimestamps Use events with kernel timestamp support (0 or 1)
                    		 ✔️ ❌

                    emu_benchmark

                    Emulation Benchmark is a set of tests aimed at measuring performance of emulated math operations performed in kernels.

                    Test name Description Params L0 OCL
                    Int64Div enqueues kernel performing an int64 division emulation
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️

                    eu_benchmark

                    EU Benchmark is a set of tests aimed at measuring performance of calculations performed in kernels.

                    Test name Description Params L0 OCL
                    DoMathOperation enqueues kernel performing a math operation. Min, Max, Modulo, And, Or and Xor are each followed by an addition, a subtraction for Max, that makes the next operand depend on the result, so their results include that operation. Integer Inc and Dec are not implemented
                    • --mixGrf Run kernels with mixed grf modes (0 or 1)
                    • --op Math operation to perform (Add or Sub or Div or Modulo or Inc or Dec or Min or Max or And or Or or Xor)
                    • --type Data type of the atomic. Keep in mind not all operations are supported for floating points (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgc Work group count
                    • --wgs Work group size
                    		 ✔️ ✔️
                    ReadAfterAtomicWrite enqueues kernel, which writes to global memory using atomic and then reads non atomically
                    • --atomic If true, write to global memory will be atomic. (0 or 1)
                    • --shuffleRead If true, each thread will write and read different memory cell. Otherwise it will be the same one. (0 or 1)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --wgs Workgroup size
                    		 ✔️ ✔️

                    gpu_cmds_benchmark

                    Gpu Commands Benchmark is a set of tests aimed at measuring GPU-side execution duration of various commands.

                    Test name Description Params L0 OCL
                    AggregatedEventSignaling measures time required to signal an aggregated event, where each operation generates an atomic write
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌
                    AggregatedEventWait measures time required to service an aggregated event wait using a single wait on a GPU
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌
                    BarrierBetweenKernels measures time required to run a barrier command between 2 kernels, including potential cache flush commands
                    • --barrierCount total amount of measured barriers
                    • --bytes bytes to flush from L3
                    • --memoryType memory type cached in L3 (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --onlyReads only reads cached in L3
                    • --remoteAccess access cached from remote tile
                    		 ✔️ ❌
                    CopyWithEvent measures time required to run a copy kernel with various event configurations.
                    • --devWaitEvent Use ZE_EVENT_SCOPE_FLAG_DEVICE for ze_event_desc_t::wait (0 or 1)
                    • --hostSignalEvent Use ZE_EVENT_POOL_HOST_VISIBLE for ze_event_pool_desc_t::flags, and use ZE_EVENT_SCOPE_FLAG_HOST for ze_event_desc_t::signal (0 or 1)
                    • --measuredCmds Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    • --timestampEvent Use ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP for ze_event_pool_desc_t::flags (0 or 1)
                    		 ✔️ ❌
                    EmptyKernel measures time required to run an empty kernel on GPU.
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    • --wgc Workgroup count
                    • --wgs Workgroup size (aka local work size)
                    		 ✔️ ❌
                    EventCtxtSwitchLatency measures context switching latency time required to switch between various engine types
                    • --firstEngine first engine to measure context switch latency (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    • --secondEngine second engine to measure context switch latency (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    		 ✔️ ❌
                    KernelWithEvent measures time required to run an empty kernel with various event configurations.
                    • --devWaitEvent Use ZE_EVENT_SCOPE_FLAG_DEVICE for ze_event_desc_t::wait (0 or 1)
                    • --hostSignalEvent Use ZE_EVENT_POOL_HOST_VISIBLE for ze_event_pool_desc_t::flags, and use ZE_EVENT_SCOPE_FLAG_HOST for ze_event_desc_t::signal (0 or 1)
                    • --inOrderCmdList Use ZE_COMMAND_LIST_FLAG_IN_ORDER for ze_command_list_desc_t::flags (0 or 1)
                    • --measuredCmds Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    • --timestampEvent Use ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP for ze_event_pool_desc_t::flags (0 or 1)
                    • --useCbe Use counter based events (0 or 1)
                    • --wgc Workgroup count
                    • --wgs Workgroup size (aka local work size)
                    		 ✔️ ❌
                    KernelWithWork measures time required to run a GPU kernel which assigns values to elements of a buffer.
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    • --usedIds Which of the get_global_id() and get_local_id() calls will be used in the kernel (None or Global or Local or AtomicPerWkg)
                    • --wgc Workgroup count
                    • --wgs Workgroup size (aka local work size)
                    		 ✔️ ❌
                    LastEventLatency Measures different ways to retrieve event from last submission.
                    • --signalOnBarrier Place signal event on barrier instead of launched kernel. (0 or 1)
                    • --useSameCmdList Synchronize on barrier using same cmd list. (0 or 1)
                    		 ✔️ ❌
                    WaitOnEventCold measures time required to service a signalled semaphore, that has never been waited for.
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌
                    WaitOnEventFromWalker measures time required to service a signalled semaphore coming from Walker command
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌
                    WaitOnEventHot measures time required to service a signalled semaphore, that was previously used
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌
                    WriteTimestamp measures time required to write a timestamp on GPU.
                    • --measuredCommands Number of commands being measured. Result is later divided by this number, to achieve time of a single command
                    		 ✔️ ❌

                    graph_api_benchmark

                    Graph Api Overhead Benchmark is a set of tests aimed at measuring CPU-side execution duration of SYCL Graphs API calls.

                    Test name Description Params L0 OCL
                    Decoder2Graph LLM decode proxy benchmark to simulate graph execution with host CPU coordination.
                    • --EmulateGraphs Emulate graph execution (1) or not (0) for L0 backend. Must be 0 for other backends (0 or 1)
                    • --NumTokens Number of tokens
                    • --UseGraphs Use graphs (1) or execute eagerly (0) (0 or 1)
                    • --UseHostTasks Use host task for single graph execution per token (1) or no host tasks for single graph per layer (0) (0 or 1)
                    • --UseNativeRecording Enables property::graph::enable_native_recording on the command graph (only valid for SYCL). With UseHostTasks=1, uses zeCommandListAppendHostFunction instead of SYCL host_task, which is unsupported in native recording mode. (0 or 1)
                    		 ✔️ ❌
                    GraphReplayLatency records a graph and appends it replayCount times to an in-order immediate command list, each time with a profiling event. Measures time from the end of one graph execution till the start of the next one.
                    • --forkCount Number of forks in the graph, each fork launches a single kernel
                    • --replayCount Number of times the graph is appended in a single iteration, must be at least 2
                    		 ✔️ ❌
                    GraphWithTransfers measures execution time of a graph captured on an immediate command list, consisting of a host-to-device copy, a kernel processing the device buffer and a device-to-host copy. Optionally, a second graph of 3 compute-bound kernels is executed in parallel on a separate immediate command list, competing for the same compute resources, and the measured time lasts until both graphs complete. Results are verified after benchmarking.
                    • --computeIterations Number of loop iterations in each compute-bound kernel of the compute graph
                    • --size Size of the host and device buffers transferred and processed by the graph
                    • --withComputeGraph Execute a second graph of 3 compute-bound kernels on a separate immediate command list in parallel with the transfer graph (0 or 1)
                    • --withCopyOffload Create the capturing immediate command list with the copy offload hint (0 or 1)
                    		 ✔️ ❌
                    MutateGraph The benchmark quantifies the benefits of using Mutable Command List for graphs compared to recreating the graphs for scratch every time
                    • --canUpdate If true, the benchmark graph changes are done using L0 mutable command list features, otherwise they are created every time from scratch. (0 or 1)
                    • --changeRate Rate at which kernels should be changed (for example, 5 means that every 5th kernel in the graph will be changed)
                    • --mutateIsa If true, the benchmark will mutate kernel ISA. (0 or 1)
                    • --numKernels Number of kernels to use in a graph
                    • --operationType Determines operation type that should be measured: initialization, modification, execution or creation (Initialize or Mutate or Execute or Create)
                    • --useInOrder If true, lists are created with ZE_COMMAND_LIST_FLAG_IN_ORDER flag, otherwise they are created without it (0 or 1)
                    		 ✔️ ❌
                    SinKernelGraph Benchmark running memory copy and kernel runs, with graphs and without graphs
                    • --UseNativeRecording Enables property::graph::enable_native_recording on the command graph (only valid for SYCL) (0 or 1)
                    • --immediateAppendCmdList Use zeCommandListImmediateAppendCommandListsExp to submit graph (only valid for L0) (0 or 1)
                    • --numKernels Number of kernel invocations
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    • --withGraphs Runs with or without graphs (0 or 1)
                    		 ✔️ ❌
                    SubmitGraph measures time spent in submitting a graph to a SYCL (or SYCL-like) queue on CPU.
                    • --EmulateGraphs Emulate graph execution (1) or use graph API (0) (0 or 1)
                    • --InOrderQueue Create the queue with the in_order property (0 or 1)
                    • --KernelExecutionTime Approximately how long a single kernel executes, in us
                    • --MeasureCompletionTime Measures time taken to complete the submission (default is to measure only submit calls) (0 or 1)
                    • --NumKernels Number of kernels to submit to the queue
                    • --Profiling Create the queue with the enable_profiling property (0 or 1)
                    • --UseEvents Do not create and track events for synchronization (0 or 1)
                    • --UseExplicit Using Explicit Graph Creation mode vs. Record and Replay (0 or 1)
                    • --UseHostTasks Submit SYCL host task after kernel enqueue (0 or 1)
                    • --UseNativeRecording Enables property::graph::enable_native_recording on the command graph (only valid for SYCL) (0 or 1)
                    		 ✔️ ✔️

                    host_function_benchmark

                    Host Function Benchmark is a set of tests aimed at measuring performance of host function execution.

                    Test name Description Params L0 OCL
                    HostFunctionCmdListImmediate Measures the overhead of the host function.
                    • --CallsCount The number of calls that are being measured.
                    • --KernelExecTime Execution time of a single kernel (in us).
                    • --MeasureCompletion Add completion time in the measurement. (0 or 1)
                    • --UseEmptyHostFunction When enabled, the test will use an empty host function, otherwise it will use a function with a 1ms busy spin. (0 or 1)
                    • --UseKernels Use kernels and a host function in between during the test. (0 or 1)
                    		 ✔️ ❌
                    HostFunctionCmdListRegular Measures the overhead of the host function.
                    • --CallsCount The number of calls that are being measured.
                    • --KernelExecTime Execution time of a single kernel (in us).
                    • --MeasureCompletion Add completion time in the measurement. (0 or 1)
                    • --UseEmptyHostFunction When enabled, the test will use an empty host function, otherwise it will use a function with a 1ms busy spin. (0 or 1)
                    • --UseKernels Use kernels and a host function in between during the test. (0 or 1)
                    		 ✔️ ❌

                    memory_benchmark

                    Memory Benchmark is a set of tests aimed at measuring bandwidth of memory transfers.

                    Test name Description Params L0 OCL
                    CopyBuffer allocates two OpenCL buffers and measures copy bandwidth between them. Buffers will be placed in device memory, if it's available.
                    • --compressedDestination Select if the destination buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --compressedSource Select if the source buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffers (Zeros or Random)
                    • --size Size of the buffers
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    CopyBufferRect allocates two OpenCL buffers and measures rectangle copy bandwidth between them. Buffers will be placed in device memory, if it's available.
                    • --dstCompressed Select if the destination buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --origin Origin of the rectangle
                    • --rPitch Row pitch of the rectangle
                    • --region Size of the rectangle
                    • --sPitch Silice pitch of the rectangle
                    • --size Size of the buffer
                    • --srcCompressed Select if the source buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    		 ❌ ✔️
                    CopyBufferToImage allocates buffer and image and measures copy bandwidth between them using immediate command list for Level Zero and command queue for OpenCL.
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --region Size of the destination image region
                    • --size Size of the buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    CopyEntireImage allocates two image objects and measures copy bandwidth between them. Images will be placed in device memory, if it's available.
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --size Size of the image
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    CopyImageRegion allocates two image objects and measures region copy bandwidth between them using immediate command list for Level Zero and command queue for OpenCL.
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --size Size of the image
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    CopyImageToBuffer allocates image and buffer and measures copy bandwidth between them using immediate command list for Level Zero and command queue for OpenCL.
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --region Size of the source image region
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    FillBuffer allocates an OpenCL buffer and measures fill bandwidth. Buffer will be placed in device memory, if it's available.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --patternSize Size of the fill pattern
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    FullRemoteAccessMemory Uses stream memory in a fashion described by 'type' to measure bandwidth of full remote memory access.
                    • --blockAccess Block access (1) or scatter access (0) (0 or 1)
                    • --elementSize Size of the single element to read in bytes (1, 2, 4, 8)
                    • --size Size of the memory to stream. Must be divisible by element size and a power of 2
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --workItems Number of work items equal to SIMD size * used hwthreads. Must be a power of 2
                    		 ❌ ✔️
                    FullRemoteAccessMemoryXeCoresDistributed Uses stream memory in a fashion described by 'type' to measure bandwidth of full remote memory accesswhen hwthreads are distributed between XeCores.
                    • --blockAccess Block access (1) or scatter access (0) (0 or 1)
                    • --elementSize Size of the single element to read in bytes (1, 2, 4, 8)
                    • --size Size of the memory to stream. Must be a power of 2
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --workItems Number of work items equal to SIMD size * used hwthreads
                    		 ❌ ✔️
                    HostMappedMemory maps a device allocation into the host address space with zeIntelMemMapDeviceMemToHost and measures the bandwidth of a CPU memcpy through that mapping. On a device with local memory the mapping is reached over the PCIe BAR and is write-combining, so writes stream while reads become uncached per-cacheline round trips, which makes the two directions differ by orders of magnitude. On an integrated device the driver returns the existing host pointer and the result is an ordinary system memory measurement. Requires a driver exposing the mapping extension and an allocation the host can alias directly.
                    • --fenceWrites Order writes with a full fence before the measurement ends, as a producer signalling a remote reader has to (0 or 1)
                    • --size Size of the mapped device allocation
                    • --type Access the CPU performs through the mapping, read or write (Read or Write or Scale or Triad or ReadConstant)
                    		 ✔️ ❌
                    MapBuffer allocates an OpenCL buffer and measures map bandwidth. Mapping operation means memory transfer from GPU to CPU or a no-op, depending on map flags.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --mapFlags OpenCL map flags passed during memory mapping (Read or Write or WriteInvalidate)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    QueueInOrderMemcpy measures time on CPU spent for multiple in order memcpy.
                    • --IsCopyOnly If true, Copy Engine is selected. If false, Compute Engine is selected (0 or 1)
                    • --count Number of memcpy operations
                    • --destinationPlacement Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --size Size of memory allocation
                    • --sourcePlacement Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    RandomAccessMemory Measures device-memory random access bandwidth for different allocation sizes, alignments and access modes.The benchmark uses 10 million accesses to memory.
                    • --accessMode Access mode to be used('Read', 'Write', 'ReadWrite')
                    • --alignment Alignment request for the allocated memory
                    • --allocationSize Size of device memory to be allocated.(Maximum supported is 16GB)
                    • --randomAccessRange Percentage of allocation size to be used for random access
                    • --randomAccessSeed Seed for generating random indexes
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    RandomAccessMultiResource Measures random access bandwidth for different allocation sizes and placements.The benchmark uses 10 million accesses to memory for each resource.
                    • --firstPlacement Placement of first resource (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --firstSize Size of first resource. (Maximum supported is 16GB)
                    • --secondPlacement Placement of second resource (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --secondSize Size of second resource. (Maximum supported is 16GB)
                    		 ✔️ ❌
                    ReadBuffer allocates an OpenCL buffer and measures read bandwidth. Read operation means transfer from GPU to CPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --reuse How hostptr allocation can be reused due to previous operations (Aligned4KB or Misaligned or Usm or Map)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    ReadBufferMisaligned allocates an OpenCL buffer and measures read bandwidth. Read operation means transfer from GPU to CPU. Destination pointer passed by the application will be misaligned by the specified amount of bytes.
                    • --misalignment Number of bytes by which misaligned the destination pointer will be misaligned
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    ReadBufferRect allocates an OpenCL buffer and measures rectangle read bandwidth. Rectangle read operation means transfer from GPU to CPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --origin Origin of the rectangle
                    • --rPitch Row pitch of the rectangle
                    • --region Size of the rectangle
                    • --sPitch Silice pitch of the rectangle
                    • --size Size of the buffer
                    		 ❌ ✔️
                    ReadDeviceMemBuffer allocates two OpenCL buffers and measures source buffer read bandwidth. Source buffer resides in device memory.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --size Size of the buffer
                    		 ❌ ✔️
                    ReadImage Measures time spent during Read Image calls
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --ptrAlignment host pointer alignment
                    • --ptrPlacement memory placement of host_ptr passed to ReadImage call (Aligned4KB or Misaligned or Usm or Map)
                    • --size Size of the image
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    RemoteAccessMemory Uses stream memory in a fashion described by 'type' to measure bandwidth with differentpercentages of remote memory access. Triad means two buffers are read and one is written to.In read and write memory is only read or written to.
                    • --remoteFraction Fraction of remote memory access. 1 / n
                    • --size Size of the memory to stream. Must be divisible by datatype size.
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --workItemSize Number of work items group together for remote check
                    		 ❌ ✔️
                    RemoteAccessMemoryMaxSaturation Uses stream memory write to measure max data bus saturation with different percentages of remote memory access
                    • --remoteFraction Fraction of remote memory access. 1 / n
                    • --size Size of the memory to stream. Must be divisible by datatype size.
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --workItemSize Number of work items group together for remote check
                    • --writesPerWorkgroup Number of work items per workgroup that access memory
                    		 ❌ ✔️
                    SLM_DataAccessLatency generates SLM local memory transactions inside thread group to measure latency between reads (uses Intel only private intel_get_cycle_counter() )
                    • --direction write or read mode (0 or 1)
                    • --occupancyDiv H/W load divider by 8, 4, 2, full occupancy
                    • --size SLM Size
                    		 ❌ ✔️
                    SlmSwitchLatency Enqueues 2 kernels with different SLM size. Measures switch time between these kernels.
                    • --firstSlmSize Size of the shared local memory per thread group. First kernel.
                    • --secondSlmSize Size of the shared local memory per thread group. Second kernel.
                    • --wgs Size of the work group.
                    		 ✔️ ❌
                    StreamAfterTransfer Goal of this test is to measure how stream kernels perform right after host to device transfer populating the data. Test does clean caches, then emits transfers and then follows with stream kernel and measures GPU execution time of it.
                    • --size Size of the memory to stream. Must be divisible by datatype size.
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    		 ❌ ✔️
                    StreamMemory Streams memory inside of kernel in a fashion described by 'type'. Copy means one memory location is read from and the second one is written to. Triad means two buffers are read and one is written to. In read and write memory is only read or written to. With --verify, Read also stores one element per 256 work groups (--verificationMode=light) or every element (--verificationMode=complete) so the output can be checked. ReadConstant is Read in which every work group reads the elements of the first work group, so all work groups read the same 'lws' elements.
                    • --contents Buffer contents zeros/random (Zeros or Random)
                    • --lws local work size
                    • --memory Memory type used for stream (Device or Host or Shared)
                    • --multiplier multiplies id used for accessing the resources to simulate partials
                    • --size Size of the memory to stream. Rounded down to a multiple of datatype size and work-group size.
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --vectorSize size of uint vector type 1/2/4/8/16
                    		 ✔️ ✔️
                    StreamMemoryImmediate Streams memory inside of kernel in a fashion described by 'type' using immediate command list. Copy means one memory location is read from and the second one is written to. Triad means two buffers are read and one is written to. In read and write memory is only read or written to. With --verify, Read also stores one element per 256 work groups (--verificationMode=light) or every element (--verificationMode=complete) so the output can be checked.
                    • --size Size of the memory to stream. Rounded down to a multiple of datatype size and work-group size.
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UnmapBuffer allocates an OpenCL buffer and measures unmap bandwidth. Unmapping operation meansmemory transfer from CPU to GPU or a no-op, depending on map flags.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --mapFlags OpenCL map flags passed during memory mapping (Read or Write or WriteInvalidate)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    UsmConcurrentCopy allocates four unified shared memory buffers, 2 in device memory and 2 in host memory. Measures concurrent copy bandwidth between them.
                    • --d2hEngine Engine used for device to host copy (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --h2dEngine Engine used for host to device copy (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of the buffer
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    UsmCopy allocates two unified shared memory buffers and measures copy bandwidth between them.
                    • --contents Contents of the buffers (Zeros or Random)
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --reuseCmdList Command list is reused between iterations (0 or 1)
                    • --size Size of the buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ✔️
                    UsmCopyConcurrentMultipleBlits Measures Copy bandwidth while performing concurrent copies between host and device using different copy engines. Engines for Host to Device copies could be selected using h2dBlitters. Engines for Device to Host copies could be selected using d2hBlitters.
                    • --d2hBlitters A bit mask for selecting copy engines to be used for device to host copy
                    • --h2dBlitters A bit mask for selecting copy engines to be used for host to device copy
                    • --size Size of the copy to be done for each copy engine
                    		 ✔️ ❌
                    UsmCopyImmediate allocates two unified shared memory buffers and measures copy bandwidth between them using immediate command list.
                    • --contents Contents of the buffers (Zeros or Random)
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --size Size of the buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    UsmCopyImmediateAfterKernels Measures CPU time of appending memory copies to an in-order immediate command list that already holds time-consuming kernels. 'Append' covers the copy appends only, including any wait for the preceding kernels to complete. 'Total' covers the time from the first kernel append until synchronization returns.
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --kernelTime Approximately how long a single kernel executes, in us
                    • --numCopies Number of zeCommandListAppendMemoryCopy calls appended after the kernels
                    • --numKernels Number of eat_time kernels appended before the copies
                    • --size Size of the copied buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    		 ✔️ ❌
                    UsmCopyMultipleBlits allocates two unified shared memory buffers, divides them into chunks, copies each chunk using a different copy engine and measures bandwidth. Results for each individual blitter engine is measured using GPU-based timings and reported separately. Total bandwidths are calculated by dividing the total buffer size by the worst result from all engines. Division of work among blitters is not always even - if main copy engine is specified (rightmost bit in --bliters argument), it gets a half of the buffer and the rest is divided between remaining copy engines. Otherwise the division is even.
                    • --blitters A bit mask for selecting copy engines
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --size Size of the operation processed by each engine
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    		 ✔️ ✔️
                    UsmCopyRegion allocates two unified shared memory buffers and measures region copy bandwidth between them using immediate command list.
                    • --contents Contents of the buffers (Zeros or Random)
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --origin Origin of the region
                    • --region Size of the region
                    • --size Size of the buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmCopyStagingBuffers Measures copy time from device/host to host/device. Host memory is non-USM allocation.Copy is done through staging USM buffers. Non-USM host ptr is never passed to L0 API, only through staging buffers.
                    • --chunks How much memory chunks should the buffer be splitted into
                    • --dst Memory placement of destination (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --size Size of the buffer
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    UsmFill allocates a unified memory buffer and measures fill bandwidth
                    • --contents Contents of the buffer (Zeros or Random)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of the buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --patternContents Select contents of the fill pattern (Zeros or Random)
                    • --patternSize Size of the fill pattern
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    UsmFillImmediate allocates a unified memory buffer and measures fill bandwidth using immediate command list
                    • --contents Contents of the buffer (Zeros or Random)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of the buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --patternContents Select contents of the fill pattern (Zeros or Random)
                    • --patternSize Size of the fill pattern
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmFillMultipleBlits allocates a unified shared memory buffer, divides it into chunks, copies each chunk using a different copy engine and measures bandwidth. Refer to UsmCopyMultipleBlits for more details.
                    • --blitters A bit mask for selecting copy engines
                    • --memory Placement of buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --patternContents Select contents of the fill pattern (Zeros or Random)
                    • --patternSize Size of the fill pattern
                    • --size Size of the operation processed by each engine
                    		 ✔️ ✔️
                    UsmFillSpecificPattern allocates a unified memory buffer and measures fill bandwidth. Allow specifying arbitrary pattern.
                    • --contents Contents of the buffer (Zeros or Random)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of the buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --pattern The fill pattern represented hexadecimally, e.g. 0x91ABCD1254
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    UsmImmediateCopyMultipleBlits allocates two unified shared memory buffers, divides them into chunks, copies each chunk using a different copy engine with an immediate command list and measures bandwidth. Results for each individual blitter engine is measured using GPU-based timings and reported separately. Total bandwidths are calculated by dividing the total buffer size by the worst result from all engines. Division of work among blitters is not always even - if main copy engine is specified (rightmost bit in --bliters argument), it gets a half of the buffer and the rest is divided between remaining copy engines. Otherwise the division is even.
                    • --blitters A bit mask for selecting copy engines
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --size Size of the operation processed by each engine
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    		 ✔️ ❌
                    UsmMemset allocates a unified memory buffer and measures memset bandwidth
                    • --contents Contents of the buffer (Zeros or Random)
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of the buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    WriteBuffer allocates an OpenCL buffer and measures write bandwidth. Write operation means transfer from CPU to GPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --reuse How hostptr allocation can be reused due to previous operations (Aligned4KB or Misaligned or Usm or Map)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    WriteBufferRect allocates an OpenCL buffer and measures rectangle write bandwidth. Rectangle write operation means transfer from CPU to GPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --contents Contents of the buffer (Zeros or Random)
                    • --inOrderQueue If set use IOQ, otherwise OOQ. Applicable only for OCL. (0 or 1)
                    • --origin Origin of the rectangle
                    • --rPitch Row pitch of the rectangle
                    • --region Size of the rectangle
                    • --reuse How hostptr allocation can be reused due to previous operations (Aligned4KB or Misaligned or Usm or Map)
                    • --sPitch Silice pitch of the rectangle
                    • --size Size of the buffer
                    		 ❌ ✔️
                    WriteImage Measures time spent during Write Image calls
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --ptrPlacement memory placement of host_ptr passed to WriteImage call (Aligned4KB or Misaligned or Usm or Map)
                    • --size Size of the image
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️

                    miscellaneous_benchmark

                    Miscellaneous Benchmark is a set of tests measuring different simple compute scenarios.

                    Test name Description Params L0 OCL
                    IoqKernelSwitchLatency measures time from end of one kernel till start of next kernel for in order queue
                    • --kernelCount Count of kernels
                    • --useEvents Use events to synchronize between kernels (0 or 1)
                    		 ❌ ✔️
                    KernelWithWork measures time required to run a GPU kernel which assigns constant values to elements of a buffer. Each thread assigns one value. Benchmark checks the impact of kernel split.
                    • --split How many times kernel is split
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --usedIds Which of the get_global_id() and get_local_id() calls will be used in the kernel (None or Global or Local or AtomicPerWkg)
                    • --wgc Workgroup count
                    • --wgs Workgroup size (aka local work size)
                    		 ❌ ✔️
                    Reduction Performs a reduction operation on a buffer. Each thread performs atomic_add on one shared memory location.
                    • --numberOfElements Number of elements that will be reduced
                    		 ❌ ✔️
                    Reduction2 Performs a reduction operation on a buffer. Each thread performs atomic_add on one shared memory location.
                    • --numberOfElements Number of elements that will be reduced
                    		 ❌ ✔️
                    Reduction3 Performs a reduction operation on a buffer. Each thread performs atomic_add on one shared memory location.
                    • --numberOfElements Number of elements that will be reduced
                    		 ❌ ✔️
                    Reduction4 Performs a reduction operation on a buffer. Each thread performs atomic_add on one shared memory location.
                    • --numberOfElements Number of elements that will be reduced
                    		 ❌ ✔️
                    Reduction5 Performs a reduction operation on a buffer. Each thread performs atomic_add on one shared memory location.
                    • --numberOfElements Number of elements that will be reduced
                    		 ❌ ✔️
                    VectorSum Performs vector addition
                    • --numberOfElementsX Number of elements in X dimension
                    • --numberOfElementsY Number of elements in Y dimension
                    • --numberOfElementsZ Number of elements in Z dimension
                    		 ❌ ✔️

                    multiprocess_benchmark

                    Multiprocess Benchmark is a set of tests aimed at measuring how different commands benefit for simultaneous execution.

                    Test name Description Params L0 OCL
                    KernelAndCopy enqueues kernel and copy operation with the ability to perform both tasks on different command queues.
                    • --runCopy Enqueue buffer to buffer copy during each iteration (0 or 1)
                    • --runKernel Enqueue kernel during each iteration (0 or 1)
                    • --twoQueues Enables using separate queues for both operations. Must be used with runCopy and runKernel (0 or 1)
                    • --useCopyQueue Use a specialized copy queue for the copy operation. Must be used with runCopy (0 or 1)
                    		 ❌ ✔️
                    MultiProcessCompute Creates a number of separate processes for each tile specified performing a compute workload and measures average time to complete all of them. Processes will use affinity mask to select specific sub-devices for the execution
                    • --opsPerKernel Operations performed in kernel, used to steer its execution time
                    • --processesPerTile Number of processes that will be started on each of the tiles specified
                    • --synchronize Synchronize all processes before each iteration (0 or 1)
                    • --tiles Tiles for execution (Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --workgroupsPerProcess Number of workgroups that each process will start
                    		 ✔️ ❌
                    MultiProcessComputeSharedBuffer Creates a number of separate processes for each tile specified performing a compute workload and measures average time to complete all of them. Processes will use affinity mask to select specific sub-devices for the execution. A single buffer for each tile is created by parent process. All processes executing on a given tile will share it via IPC calls.
                    • --processesPerTile Number of processes that will be started on each of the tiles specified
                    • --synchronize Synchronize all processes before each iteration (0 or 1)
                    • --tiles Tiles for execution (Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --workgroupsPerProcess Number of workgroups that each process will start
                    		 ✔️ ❌
                    MultiProcessImmediateCmdlistCompletion measures completion latency of AppendMemoryCopy issued from multiple processes to Immediate Command Lists.Engines to be used for submissions are selected based on the enabled bits of engineMask.Bits of the 'engineMask' are indexed from right to left. So rightmost bit represents first engine and leftmost, the last engine.If 'numberOfProcesses' is greater than selected engine count, then the excess processes are assigned to selected engines one each, in a round-robin method.if selected engineCount == 1, then all processes are assigned to that engine.
                    • --copySize copy size in bytes
                    • --engineGroup engine group to be used
                    • --engineMask bit mask for selecting engines to be used for submission
                    • --numberOfProcesses total number of processes
                    		 ✔️ ❌
                    MultiProcessImmediateCmdlistSubmission measures submission latency of walker command issued from multiple processes to Immediate Command Lists.If 'numberOfProcesses' is greater than engine count, then the excess processes are assigned to engines one each, in a round-robin method.if engineCount == 1, then all processes are assigned to the engine.
                    • --numberOfProcesses total numer of processes
                    		 ✔️ ❌
                    MultiProcessInit Measures the initialization overhead in a multi-process application.For Level Zero we only measure the first invocation of zeInit() per process execution.
                    • --initFlag Initialization flag. For Level Zero: 0 - default, 1 - ZE_INIT_FLAG_GPU_ONLY, 2 - ZE_INIT_FLAG_VPU_ONLY
                    • --numberOfProcesses Total number of processes
                    		 ✔️ ❌

                    multithread_benchmark

                    Multithread Benchmark is a set of tests aimed at measuring how different commands benefit from multithreaded execution.

                    Test name Description Params L0 OCL
                    ImmediateCommandListCompletion measures completion latency of AppendMemoryCopy issued from multiple threads to Immediate Command Lists.Engines to be used for submissions are selected based on the enabled bits of engineMask.'threadsPerEngine' number of threads submits commands to each selected engine.If 'numberOfThreads' is greater than 'threadsPerEngine' x selected engine count, then the excess threads are assigned to selected engines one each, in a round-robin method.if selected engineCount == 1, then all threads are assigned to that engine.
                    • --copySize copy size in bytes
                    • --engineGroup engine group to be used
                    • --engineMask bit mask for selecting engines to be used for submission
                    • --numberOfThreads total number of threads
                    • --threadsPerEngine number of threads submitting commands to each engine
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    ImmediateCommandListSubmission measures submission latency of AppendLaunchKernel issued from multiple threads to Immediate Command Lists.'threadsPerEngine' count of threads submit commands to each engine.If 'numberOfThreads' is greater than 'threadsPerEngine' x engine count, then the excess threads are assigned to engines one each, in a round-robin method.if engineCount == 1, then all threads are assigned to the engine.
                    • --numberOfThreads total number of threads
                    • --threadsPerEngine number of threads submitting commands to each engine
                    		 ✔️ ❌
                    SvmCopy enqueues multiple svm copies on multiple threads concurrently.
                    • --numberOfThreads Number of threads that will run concurrently
                    		 ✔️ ✔️

                    multitile_memory_benchmark

                    Multi-tile Memory Benchmark is a set of tests aimed at measuring bandwidth of memory transfers performed on a multi-tile device.

                    Test name Description Params L0 OCL
                    CopyBuffer allocates two OpenCL buffers and measures copy bandwidth between them. Buffers will be placed in device memory, if it's available.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --dst Placement of memory for the destination buffer (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --dstCompressed Select if the destination buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffers
                    • --src Placement of memory for the source buffer (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --srcCompressed Select if the source buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    FillBuffer allocates an OpenCL buffer and measures fill bandwidth. Buffer will be placed in device memory, if it's available.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of memory for the buffer (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --patternSize Size of the fill pattern
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    ReadBuffer allocates an OpenCL buffer and measures read bandwidth. Read operation means transfer from GPU to CPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --memory Placement of memory for the buffer (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️
                    UsmBidirectionalCopy allocates two unified device memory buffers, each on a different tile, and measures copy bandwidth between. Test measures copies on two directions, which can be controlled with the -write parameter: with -write=1, each tile performs a write operation. For instance: queue is placed in tile 0, source is buffer in tile 0, and destination is in tile 1. Similarly for tile 1, queue is placed in tile 1, source in tile 1, and destination in tile 0. With -write=0, the destination and source are flipped: queue is placed in tile 0, source is buffer in tile 1, and destination is in tile 0, while for tile 1, queue is placed in tile 1, source in tile 0, and destination in tile 1.
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --size Size of the buffers
                    • --write Which operation is used, whether write or read (0 or 1)
                    		 ✔️ ❌
                    UsmCopy allocates two unified shared memory buffers and measures copy bandwidth between them using a builtin function.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --dst Placement of memory for the destination buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffers
                    • --src Placement of memory for the source buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    UsmCopyImmediate allocates two unified shared memory buffers and measures copy bandwidth between them using a builtin function appended to an immediate list.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --dst Placement of memory for the destination buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffers
                    • --src Placement of memory for the source buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)
                    		 ✔️ ❌
                    UsmCopyKernel allocates two unified shared memory buffers and measures copy bandwidth between them using a custom kernel.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --dst Placement of memory for the destination buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffers
                    • --src Placement of memory for the source buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    UsmFill allocates a unified shared memory buffer and measures fill bandwidth.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --memory Placement of memory for the buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --patternSize Size of the fill pattern
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ✔️
                    UsmSharedMigrateCpu allocates a unified shared memory buffer and measures time to migrate it from GPU to CPU.
                    • --accessAllBytes Select, whether entire resource or only one byte will be accessed on CPU (0 or 1)
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --memory Placement of memory for the buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --size Size of the buffer
                    		 ✔️ ✔️
                    UsmSharedMigrateGpu allocates a unified shared memory buffer and measures time to migrate it from CPU to GPU.
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --memory Placement of memory for the buffer (Host or Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --size Size of the buffer
                    		 ✔️ ✔️
                    WriteBuffer allocates an OpenCL buffer and measures write bandwidth. Write operation means transfer from CPU to GPU.
                    • --compressed Select if the buffer is to be compressed. Will be skipped, if device does not support compression (0 or 1)
                    • --context How context will be created (Root or Tile0 or Tile1 or Tile2 or Tile3 or a list separated with ':')
                    • --memory Placement of memory for the buffer (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --queue Which device within the context will perform the operation (Root or Tile0 or Tile1 or Tile2 or Tile3)
                    • --size Size of the buffer
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ❌ ✔️

                    p2p_benchmark

                    P2P Benchmark is a set of tests aimed at measuring bandwidth and latency of memory transfers between peer devices.

                    Test name Description Params L0 OCL
                    UsmCopyBidirectionalRead allocates unified device memory buffers on two root devices and measures the bandwidth of two copies at the same time, each device reading from the other on an immediate command list. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --localDeviceId Local root device
                    • --localEngine Engine of the local device, which reads from the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --peerDeviceId Peer root device
                    • --peerEngine Engine of the peer device, which reads from the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmCopyBidirectionalWrite allocates unified device memory buffers on two root devices and measures the bandwidth of two copies at the same time, each device writing to the other on an immediate command list. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --localDeviceId Local root device
                    • --localEngine Engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --peerDeviceId Peer root device
                    • --peerEngine Engine of the peer device, which writes to the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmCopyMultipleBlits allocates two unified device memory buffers on separate devices and performs a copy between sections (or chunks) of these using a different copy engine and measures bandwidth. Test first checks for P2P capabilities in the target platform before submitting the copy. Results for each individual blitter engine is measured using GPU-based timings and reported separately. Total bandwidths are calculated by dividing the total buffer size by the worst result from all engines. Division of work among blitters is not always even - if main copy engine is specified (rightmost bit in --bliters argument), it gets a half of the buffer and the rest is divided between remaining copy engines. Otherwise the division is even.
                    • --blitters A bit mask for selecting copy engines
                    • --dstDeviceId Destination device
                    • --size Size of the operation processed by each engine
                    • --srcDeviceId Source device
                    		 ✔️ ❌
                    UsmCopyPushPull allocates unified device memory buffers on two root devices and measures the bandwidth of two copies at the same time on immediate command lists: the local device writes to the peer while the peer reads from the local device, so both copies move data in the same direction. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --readEngine Engine of the peer device, which reads from the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --writeEngine Engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    		 ✔️ ❌
                    UsmCopyRead allocates unified device memory buffers on two root devices and measures the bandwidth of a copy from the peer device to the local device, submitted by the local device on an immediate command list, so the local device reads from the peer.
                    • --engine Engine of the local device, which reads from the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmCopyWrite allocates unified device memory buffers on two root devices and measures the bandwidth of a copy from the local device to the peer device, submitted by the local device on an immediate command list, so the local device writes to the peer.
                    • --engine Engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmCopyWriteRead allocates unified device memory buffers on two root devices and measures the bandwidth of two copies at the same time, both submitted by the local device on two engines with immediate command lists: one writes to the peer and one reads from the peer, so the link carries data in both directions. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --readEngine Engine of the local device that reads from the peer. Must differ from writeEngine (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --writeEngine Engine of the local device that writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    		 ✔️ ❌
                    UsmEUCopy allocates two unified device memory buffers on separate devices, performs a copy between them using a compute engine, and reports bandwidth. Test first checks for P2P capabilities in the target platform before submitting the copy.
                    • --contents Contents of the buffers (Zeros or Random)
                    • --dstDeviceId Destination device
                    • --reuseCmdList Command list is reused between iterations (0 or 1)
                    • --size Size of the buffer
                    • --srcDeviceId Source device
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEUReduceCopyRing connects numDevices root devices in a ring and runs on every device the kernels of a large message allreduce ring with EU writes to the peer. The message is split in numDevices parts and each part is moved in chunks. For every chunk a copy kernel writes a part to a temporary buffer on the next device, a barrier kernel waits for the previous device, and a reduce-copy kernel computes out1 = in1 + in2 in local device memory and writes out1 to the next device. With more than 2 devices the barrier and the reduce-copy repeat for every ring step and copy kernels for the allgather follow. The kernels are compiled at run time from p2p_benchmark_reduce_copy.cl, which can be modified next to the binary without rebuilding. All devices start together, so both directions of every link carry writes at the same time. Reports the bandwidth of the peer writes of one device, which is the allreduce bus bandwidth. With throttledWorkItems one reduce-copy kernel over the whole message runs with rate limited peer writes, which avoids the collapse of two-way EU peer writes and shows the peak bandwidth of the link. The best count is just below the collapse and depends on the system.
                    • --numDevices Number of root devices in the ring. Test is skipped when the system has fewer devices
                    • --size Size of the allreduce message on every device. With throttledWorkItems it is the size reduced and written to the peer by each device
                    • --throttledWorkItems If not 0, run instead of the allreduce one throttled reduce-copy kernel from p2p_benchmark_reduce_copy_throttled.cl with this many work-items per device. The work-items loop over the message and after every iteration one thread per work-group pauses while the work-group waits on a barrier, so this count sets the rate of peer writes
                    • --tmpBufferSize Size of the temporary buffer on every device. The message is moved in chunks of tmpBufferSize / max(numDevices, 3). The default is 384MB
                    • --useEvents Report the longest time across devices from the start of the first kernel to the end of the last kernel, measured with GPU timestamps. Otherwise report the CPU time from releasing the kernels until all devices complete (0 or 1)
                    		 ✔️ ❌
                    UsmEuCopyBidirectionalWriteChunked allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes in both directions at the same time, each device writing to the other. The copy kernel from p2p_benchmark_reduce_copy.cl copies the source buffer to the peer. It is launched once per chunk of chunkSize on an in-order immediate command list of a compute engine, so the peer writes in flight are limited to about one chunk. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. With useEvents the time of a copy is measured from the start of its first chunk to the end of its last chunk. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --chunkSize Size of each chunk
                    • --localDeviceId Local root device
                    • --localEngine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --peerDeviceId Peer root device
                    • --peerEngine Compute engine of the peer device, which writes to the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuCopyBidirectionalWriteThrottled allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes in both directions at the same time, each device writing to the other. The copy kernel from p2p_benchmark_copy_throttled.cl copies the source buffer to the peer with throttledWorkItems work-items on a compute engine. The work-items loop over the buffer and after every iteration one thread per work-group pauses while the work-group waits on a barrier, so the work-item count sets the rate of the peer writes. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. The best count depends on the system. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --localDeviceId Local root device
                    • --localEngine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --peerDeviceId Peer root device
                    • --peerEngine Compute engine of the peer device, which writes to the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --throttledWorkItems Number of work-items of each copy kernel, a multiple of 64
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuCopyWriteChunked allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes from the local device to the peer device. The copy kernel from p2p_benchmark_reduce_copy.cl copies the source buffer to the peer. It is launched once per chunk of chunkSize on an in-order immediate command list of a compute engine, so the peer writes in flight are limited to about one chunk. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. With useEvents the time of a copy is measured from the start of its first chunk to the end of its last chunk.
                    • --chunkSize Size of each chunk
                    • --engine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuCopyWriteThrottled allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes from the local device to the peer device. The copy kernel from p2p_benchmark_copy_throttled.cl copies the source buffer to the peer with throttledWorkItems work-items on a compute engine. The work-items loop over the buffer and after every iteration one thread per work-group pauses while the work-group waits on a barrier, so the work-item count sets the rate of the peer writes. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. The best count depends on the system.
                    • --engine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --throttledWorkItems Number of work-items of each copy kernel, a multiple of 64
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuReduceCopyBidirectionalWriteChunked allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes in both directions at the same time, each device writing to the other. The reduce-copy kernel from p2p_benchmark_reduce_copy.cl computes out1 = in1 + in2 in local device memory and writes out1 to the peer. It is launched once per chunk of chunkSize on an in-order immediate command list of a compute engine, so the peer writes in flight are limited to about one chunk. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. With useEvents the time of a copy is measured from the start of its first chunk to the end of its last chunk. The copies are released together by one host-signalled event. Reports each copy and 'AllCopies', the bytes of both copies divided by the time until the last copy completes, which is always measured with the CPU timer.
                    • --chunkSize Size of each chunk
                    • --localDeviceId Local root device
                    • --localEngine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --peerDeviceId Peer root device
                    • --peerEngine Compute engine of the peer device, which writes to the local device (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuReduceCopyWriteChunked allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes from the local device to the peer device. The reduce-copy kernel from p2p_benchmark_reduce_copy.cl computes out1 = in1 + in2 in local device memory and writes out1 to the peer. It is launched once per chunk of chunkSize on an in-order immediate command list of a compute engine, so the peer writes in flight are limited to about one chunk. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. With useEvents the time of a copy is measured from the start of its first chunk to the end of its last chunk.
                    • --chunkSize Size of each chunk
                    • --engine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmEuReduceCopyWriteThrottled allocates unified device memory buffers on two root devices and measures the bandwidth of EU writes from the local device to the peer device. The reduce-copy kernel from p2p_benchmark_reduce_copy_throttled.cl computes out1 = in1 + in2 in local device memory and writes out1 to the peer with throttledWorkItems work-items on a compute engine. The work-items loop over the buffer and after every iteration one thread per work-group pauses while the work-group waits on a barrier, so the work-item count sets the rate of the peer writes. Limiting the data in flight avoids the bandwidth collapse of concurrent peer writes on some systems. The best count depends on the system.
                    • --engine Compute engine of the local device, which writes to the peer (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                    • --localDeviceId Local root device
                    • --peerDeviceId Peer root device
                    • --size Size of each copy
                    • --throttledWorkItems Number of work-items of each copy kernel, a multiple of 64
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    		 ✔️ ❌
                    UsmImmediateCopyMultipleBlits allocates two unified device memory buffers on separate devices and performs a copy between sections (or chunks) of these using a different copy engine with an immediate command list and measures bandwidth. Test first checks for P2P capabilities in the target platform before submitting the copy. Results for each individual blitter engine is measured using GPU-based timings and reported separately. Total bandwidths are calculated by dividing the total buffer size by the worst result from all engines. Division of work among blitters is not always even - if main copy engine is specified (rightmost bit in --bliters argument), it gets a half of the buffer and the rest is divided between remaining copy engines. Otherwise the division is even.
                    • --blitters A bit mask for selecting copy engines
                    • --dstDeviceId Destination device
                    • --size Size of the operation processed by each engine
                    • --srcDeviceId Source device
                    		 ✔️ ❌

                    page_fault_benchmark

                    Page Fault Benchmark is a set of tests aimed at measuring bandwidth of memory migration.

                    Test name Description Params L0 OCL
                    NonUsmCopy Measures time for non usm transfers and further potential host updates, updates/reallocates memory on the host between copies, uses immediate command lists and copy offload
                    • --dst Placement of the destination buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --prefetch prefetch shared system buffer before each copy (0 or 1)
                    • --reallocate reallocate buffers before each copy (0 or 1)
                    • --size Size of the buffer
                    • --src Placement of the source buffer (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --updateOnHost memory is updated on the host between copies (0 or 1)
                    		 ✔️ ❌
                    NonUsmStreamMemory Streams memory inside of kernel in a fashion described by 'type'. Copy means one memory location is read from and the second one is written to. Triad means two buffers are read and one is written to. In read and write memory is only read or written to. With --verify, Read also stores one element per 256 work groups (--verificationMode=light) or every element (--verificationMode=complete) so the output can be checked.
                    • --contents Buffer contents zeros/random (Zeros or Random)
                    • --lws local work size
                    • --memoryPlacement Memory type used for stream (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                    • --multiplier multiplies id used for accessing the resources to simulate partials
                    • --prefetch prefetch shared system buffer before each copy (0 or 1)
                    • --size Size of the memory to stream. Rounded down to a multiple of datatype size and work-group size.
                    • --type Memory streaming type (Read or Write or Scale or Triad or ReadConstant)
                    • --useEvents Perform GPU-side measurements using events (0 or 1)
                    • --vectorSize size of uint vector type 1/2/4/8/16
                    		 ✔️ ❌
                    UsmSharedMigrateCpu allocates a unified shared memory buffer and measures bandwidth for kernel that must migrate resource from GPU to CPU
                    • --accessAllBytes Select, whether entire resource or only one byte will be accessed on CPU (0 or 1)
                    • --preferredLocation Apply memadvise with preferred device location (system, device, none) (System or Device or None)
                    • --size Size of the buffer
                    		 ✔️ ✔️
                    UsmSharedMigrateGpu allocates a unified shared memory buffer and measures bandwidth for kernel that must migrate resource from CPU to GPU
                    • --preferredLocation Apply memadvise with preferred device location (system, device, none) (System or Device or None)
                    • --prefetch Explicitly migrate shared allocation to device associated with command queue (0 or 1)
                    • --size Size of the buffer
                    		 ✔️ ✔️
                    UsmSharedMigrateGpuForFill allocates a unified shared memory buffer and measures bandwidth for memory fill operation that must migrate resource from CPU to GPU
                    • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                    • --preferredLocation Apply memadvise with preferred device location (system, device, none) (System or Device or None)
                    • --prefetch Explicitly migrate shared allocation to device associated with command queue (0 or 1)
                    • --size Size of the buffer
                    		 ✔️ ✔️

                    record_and_replay_benchmark

                    Record and Replay Benchmark is a set of tests aimed at measuring performance of recording and execution of commandlist graphs.

                    Test name Description Params L0 OCL
                    RecordGraph measures overhead of recording a graph.
                    • --emulate Emulate record and replay graph API using regular commandlists. (0 or 1)
                    • --mDest Take graph destroy phase into account in perf measurment. (0 or 1)
                    • --mInst Take graph instantion phase into account in perf measurment. (0 or 1)
                    • --mRec Take graph recording phase into account in perf measurment. (0 or 1)
                    • --nAppendCopy Number of appendCopy calls per command set
                    • --nAppendKern Number of appendLaunchKernel calls per command set.
                    • --nCmdSetsInLvl Number of command sets per level.
                    • --nForksInLvl Number of forks to introduce in the graph per graph level.
                    • --nInstantiations Number of executable graphs to instantiate.
                    • --nLvls Number of levels (nesting) in the graph using forks.
                    		 ✔️ ❌

                    torch_benchmark

                    Torch Benchmark is a set of tests aimed at measuring performance of workloads generated by PyTorch.

                    Test name Description Params L0 OCL
                    KernelSubmitEventRecordQuery Measures the time to query status of completed event, multiple times.
                    • --EventQueryIterations Number of event status queries after event record.
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --Profiling Create queues with profiling enabled. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitEventRecordWait Measures the time taken to record an event for a submitted kernel, and waiting for this kernel to complete.
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --Profiling Create queues with profiling enabled. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitGraphMultiQueue Measures the performance of submission + synchronization of kernels using a captured command graph on multiple queues.
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --KernelsPerQueue Number of kernels recorded in a graph
                    • --Profiling Create the queue with the enable_profiling property (SYCL only). (0 or 1)
                    • --UseEvents Use events when enqueuing kernels. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitGraphSingleQueue Measures the performance of batches of submission + synchronization of kernels using a captured command graph on a single in-order queue.
                    • --KernelBatchSize Size of a batch of kernels after which synchronization occurs.
                    • --KernelName Name of the kernel. (Empty or Add or AddSequence)
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --KernelsPerQueue Number of kernels recorded in a graph.
                    • --Profiling Create the queue with the enable_profiling property (SYCL only). (0 or 1)
                    • --UseEvents Use events when enqueuing kernels. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitGraphVllmMock Measures the overhead of SYCL graph submission using mock kernels that represent the computation of Llama-3.2-1B model with the exception that gemm kernel is not included. The kernels do not perform any real computation but have similar launch parameters and resource usage as the actual kernels used in the model.
                    • --AllocCount Number of allocations.
                    • --GraphScenario Graph scenario to run. 0: combined graphs 1-3, 1: layer normalization, 2: self-attention layer, 3: post attention layer normalization.
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --Profiling Create the queue with the profiling enabled. (0 or 1)
                    • --UseEvents Use SYCL events for kernel submissions (true) or eventless submit via nd_launch (false). (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitLinearKernelSize Measures submit with linear kernel size.
                    • --KernelBatchSize Size of a batch of kernels after which synchronization occurs. 0 means one synchronization after all kernels are submitted.
                    • --KernelSize Size of the kernel (allowed: 32, 128, 512, 1024, 5120).
                    • --Profiling Create the queue with the enable_profiling property (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitMemoryReuse Measures submit kernel with memory reuse.
                    • --KernelBatchSize Size of a batch of kernels after which synchronization occurs. 0 means one synchronization after all kernels are submitted.
                    • --KernelDataType Data type passed to kernel. (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --Profiling Create the queue with the enable_profiling property (0 or 1)
                    • --UseEvents Use events when enqueuing kernels. When false, SYCL will use eventless enqueue functions. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitMultiQueue Measures submit kernel from 2 queues.
                    • --KernelWGCount Number of workgroups.
                    • --KernelWGSize Size of each workgroup.
                    • --KernelsPerQueue Number of kernels per queue.
                    • --MeasureCompletionTime Measures total time taken to complete all submissions and waits at the end of iteration. By default only kernel submission time is measured. (0 or 1)
                    • --Profiling Create the queue with the enable_profiling property (0 or 1)
                    • --UseEvents Use events when enqueuing kernels. (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitSingleQueue Measures time spent on executing simple math kernels.
                    • --KernelBatchSize Size of each batch of kernels before syncing. If 0, no batching is performed
                    • --KernelDataType Data type for kernel operations (Int32 or Int64 or Float or Double or CopyableObject or Mixed)
                    • --KernelName Name of the kernel to execute (Empty or Add or AddSequence)
                    • --KernelParamsNum Number of kernel parameters for add operation
                    • --KernelSubmitPattern Pattern for submitting kernels (Single or D2h_after_batch or H2d_before_batch)
                    • --KernelWGCount Number of workgroups
                    • --KernelWGSize Size of each workgroup
                    • --Profiling Create the queue with the enable_profiling property (0 or 1)
                    • --UseEvents Use events when enqueuing kernels (0 or 1)
                    		 ✔️ ❌
                    KernelSubmitSlmSize Measures submit with SLM size specified.
                    • --KernelBatchSize Sychronization interval.
                    • --MeasureCompletionTime Measures total time of one batch submission (multiple submits) together with wait after the batch. Default is to measure one submit call only. (0 or 1)
                    • --Profiling Create the queue with the enable_profiling property. Currently, only in-order queue is supported. (0 or 1)
                    • --SlmNum Size of used shared local memory.
                    		 ✔️ ❌

                    ulls_benchmark

                    Ulls Benchmark is a set of tests aimed at measuring Ultra Low Latency Submission (ULLS) performance impact.

                    Test name Description Params L0 OCL
                    BestSubmission enqueues a system memory write via PIPE_CONTROL and measures when update becomes visible on the CPU.
                      		✔️ ❌
                      BestWalkerNthCommandListSubmission enqueues single kernel on n command lists, which updates system memory location and then busy-loops on CPU until the update of the kernel of nth command list becomes visible. This is L0 only test.
                      • --CmdListCount Command list count
                      		 ✔️ ❌
                      BestWalkerNthSubmission enqueues n kernels, which updates system memory location and then busy-loops on CPU until the update of nth kernel becomes visible.
                      • --KernelCount Kernel count
                      		 ✔️ ❌
                      BestWalkerNthSubmissionImmediate enqueues n kernels, which updates system memory location and then busy-loops on CPU until the update of nth kernel becomes visible. Kernel is enqueued using low-latency immediate command list, so the test is LevelZero-specific.
                      • --KernelCount Kernel count
                      		 ✔️ ❌
                      BestWalkerSubmission enqueues kernel, which updates system memory location and then busy-loops on CPU until the update becomes visible.
                        		✔️ ✔️
                        BestWalkerSubmissionImmediate enqueues kernel, which updates system memory location and then busy-loops on CPU until the update becomes visible. Kernel is enqueued using low-latency immediate command list, so the test is LevelZero-specific.
                          		✔️ ❌
                          BestWalkerSubmissionImmediateMultiCmdlists Append N kernels on N cmdlists, which updates system memory locations and then waits using busy-loop on CPU until the update becomes visible. Kernels are appended using immediate command lists.Amount of command lists is specified by cmdlistCount.
                          • --cmdlistCount Count of command lists
                          		 ✔️ ❌
                          CompletionLatency enqueues system memory write and measures time between the moment, when update is visible on CPU and the moment, when synchronizing call returns.
                            		✔️ ❌
                            CopyPriorityPreemption submits a long memory copy on normal-priority queue, then a short copy on high-priority queue, and measures GPU time from normal-priority copy start to high-priority copy end.
                            • --dst Placement of destination memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                            • --forceBlitter Force blitter engine. Test will be skipped if device does not support blitter. Warning: in OpenCL blitter may still be used even if not forced (0 or 1)
                            • --longCopySize Size of normal-priority (long) memory copy
                            • --shortCopySize Size of high-priority (short) memory copy
                            • --src Placement of source memory (Device or Host or Shared or non-USM-mapped or non-USMmisaligned or non-USM4KBAligned or non-USM2MBAligned or non-USMmisaligned-imported or non-USM4KBAligned-imported or non-USM2MBAligned-imported or non-USM)
                            		 ✔️ ❌
                            CopySubmissionEvents enqueues 4 byte copy to copy engine and return submission delta which is time between host API call and copy engine start
                            • --engine Engine used for copying (RCS or CCSX (X = 0-3) or BCSX (X = 0-8))
                            		 ✔️ ✔️
                            EmptyKernel enqueues empty kernel and measures time to launch it and wait for it on CPU, thus measuring walker spawn time.
                            • --wgc Workgroup count
                            • --wgs Workgroup size (aka local work size)
                            		 ✔️ ✔️
                            EmptyKernelImmediate enqueues empty kernel and measures time to launch it using immediate command list and wait for it on CPU, thus measuring walker spawn time.
                            • --UseEventForHostSync If true, use events to synchronize with host.If false, use zeCommandListHostSynchronize (0 or 1)
                            • --wgc Workgroup count
                            • --wgs Workgroup size (aka local work size)
                            		 ✔️ ❌
                            EmptyKernelsWithGlobalTimer submits multiple kernels with signal events and barriers in single command lists, uses global timestamps to get total time of their execution
                            • --count Count of kernels
                            		 ✔️ ❌
                            EnqueueBarrierWithEmptyWaitlist enqueues kernel with barriers with empty waitlists inbetween, waiting on the last barriers event
                            • --enqueueCount Number of enqueues
                            • --outOfOrderQueue Use out of order queue (0 or 1)
                            		 ❌ ✔️
                            KernelSwitch measures time from end of one kernel till start of next kernel
                            • --barrier synchronization with barrier instead of events (0 or 1)
                            • --count Count of kernels
                            • --ctrBasedEvents use counter based events for in order (0 or 1)
                            • --hostVisible events are with host visible flag (0 or 1)
                            • --ioq use in order queue/command list (0 or 1)
                            • --kernelTime Approximately how long a single kernel executes, in us
                            		 ✔️ ✔️
                            KernelSwitchFill measures time from end of one kernel till start of next fill kernel
                            • --count Count of kernels
                            • --fillSize Size of the fill
                            • --ioq Use implicit dependency between kernels (0 or 1)
                            		 ✔️ ✔️
                            KernelSwitchImm measures time from end of one kernel till start of next kernel using immediate command lists
                            • --barrier synchronization with barrier instead of events (0 or 1)
                            • --count Count of kernels
                            • --cpuVisible events are with host visible flag (0 or 1)
                            • --ctrEvents use counter based events for in order (0 or 1)
                            • --ioq use in order queue/command list (0 or 1)
                            • --kTime Approximately how long a single kernel executes, in us
                            • --profiling use profiling to obtain switch time (0 or 1)
                            		 ✔️ ❌
                            KernelSwitchPriority submits kernels on queues with different priorities and measures kernel switch latency on single queue.
                            • --kernelCount How many kernels to switch between
                            • --kernelExecutionTime Time for each kernel execution
                            • --measuredQueuePriority Priority level of the queue where kernels being measured are submitted (Low or Normal or High)
                            • --secondaryQueuePriority Priority level of the other queue submitting kernels to create contention (Low or Normal or High)
                            • --useIoq Use implicit synchronization for kernel execution (0 or 1)
                            		 ✔️ ❌
                            KernelSwitchPriorityImm submits kernels on queues with different priorities and measures kernel switch latency on single queue on immediate cmd lists.
                            • --kernelCount How many kernels to switch between
                            • --kernelExecutionTime Time for each kernel execution
                            • --measuredQueuePriority Priority level of the queue where kernels being measured are submitted (Low or Normal or High)
                            • --secondaryQueuePriority Priority level of the other queue submitting kernels to create contention (Low or Normal or High)
                            • --useIoq Use implicit synchronization for kernel execution (0 or 1)
                            		 ✔️ ❌
                            KernelWithWork measures time required to run a GPU kernel which assigns constant values to elements of a buffer. Each thread assigns one value.
                            • --usedIds Which of the get_global_id() and get_local_id() calls will be used in the kernel (None or Global or Local or AtomicPerWkg)
                            • --wgc Workgroup count
                            • --wgs Workgroup size (aka local work size)
                            		 ✔️ ✔️
                            KernelWithWorkImmediate measures time required to run a GPU kernel which assigns constant values to elements of a buffer using immediate command list. Each thread assigns one value.
                            • --UseEventForHostSync If true, use events to synchronize with host.If false, use zeCommandListHostSynchronize (0 or 1)
                            • --usedIds Which of the get_global_id() and get_local_id() calls will be used in the kernel (None or Global or Local or AtomicPerWkg)
                            • --wgc Workgroup count
                            • --wgs Workgroup size (aka local work size)
                            		 ✔️ ❌
                            KernelWithWorkPeriodic measures average time required to run a GPU kernel which assigns constant values to elements of a buffer. Each thread assigns one value. Kernel is run multiple times with a set delay between submissions.
                            • --numSubmissions Number of kernel enqueues to run
                            • --timeBetweenSubmissions Delay between kernel enqueues in microseconds
                            		 ✔️ ❌
                            MultiArgumentKernelSwitchLatency measures time from end of one kernel till start of next kernel with different argument patterns
                            • --argumentCount Argument count in a kernel - supported values: 1, 4, 8, 16, 32, 64
                            • --kernelCount Count of kernels
                            • --profiling Use profiling on the gpu to collect score (0 or 1)
                            • --workgroupCount Count of workgroups
                            • --workgroupSize Size of the workgroup
                            		 ✔️ ❌
                            MultiKernelExecution submits multiple kernel in single command lists, measures total time of their execution
                            • --count Count of kernels within command list
                            • --delay how much delay between atomic reads
                            • --inOrderOverOOO use out of order queue to implement in order queue (0 or 1)
                            • --profiling use profiling on the gpu to collect score (0 or 1)
                            • --useMCL Use mutable command list (0 or 1)
                            • --wkgCount work group count of each kernel
                            • --wkgSizes work group size of each kernel
                            		 ✔️ ❌
                            MultiQueueExecution submits empty kernels on queues with different priorities and measures total execution time of both queues.
                            • --kernelCount How many kernels to switch between
                            • --measuredQueuePriority Priority level of the queue where kernels being measured are submitted (Low or Normal or High)
                            • --secondaryQueuePriority Priority level of the other queue submitting kernels to create contention (Low or Normal or High)
                            • --useIoq Use implicit synchronization for kernel execution (0 or 1)
                            • --useSecondaryQueue Use secondary cmdlist (0 or 1)
                            		 ✔️ ❌
                            MultiQueueSubmission enqueues kernel on multiple command queues
                            • --queueCount Number of command queues created
                            • --wgc Workgroup count
                            • --wgs Workgroup size
                            		 ✔️ ✔️
                            MultipleImmediateWithDependencies Creates N immediate command lists. Submits kernels in order to each of thoseEach kernel has a dependency on previous oneSubmissions are small to allows concurrent executionMeassures time from scheduling start, till all command lists are completed
                            • --UseEventForHostSync If true, use events to synchronize with host. If false, use zeCommandListHostSynchronize (0 or 1)
                            • --cmdlistCount Count of command lists
                            • --submissionsPerQueue Number of kernel submissions per command list
                            		 ✔️ ❌
                            NewResourcesSubmissionDevice enqueues kernel that uses a buffer placed in device memory to measure resource preparation time. The resource is destroyed and recreated for each iteration to ensure it is a different memory allocation.
                            • --size Size of the buffer
                            		 ✔️ ✔️
                            NewResourcesSubmissionHost enqueues kernel that uses a buffer placed in host memory to measure resource preparation time. The resource is destroyed and recreated for each iteration to ensure it is a different memory allocation.
                            • --size Size of the buffer
                            		 ✔️ ✔️
                            NewResourcesWithGpuAccess enqueues kernel that accesses an entire buffer placed in device memory to measure resource preparation time. The resource is destroyed and recreated for each iteration to ensure it is a different memory allocation.
                            • --size Size of the buffer
                            		 ✔️ ✔️
                            QueueConcurrency Submits multiple kernels to out of order queue returning events. There is a sequence of long kernel, short kernel, wait for short kernel Then calls synchronization and meassures performance
                            • --kernelCount How many kernels are submitted
                            • --kernelTime How long each work item is in kernel
                            • --workgroupCount Workgroup Count of each kernel
                            		 ❌ ✔️
                            QueuePriorities Uses queues with different priorities to meassure submission and context switch latencies
                            • --highTime How long each work item is in high priority kernel
                            • --lowTime How long each work item is in low priority kernel
                            • --priorities Low priority command queue property is used (0 or 1)
                            • --sleep sleep time in us after low priority kernel flushed
                            • --wgc Workgroup count of high priority kernel
                            		 ❌ ✔️
                            QueueSwitch creates multiple queues, creates dependencies with events between those and measures switch time
                            • --differentKernels Alternate between two different kernels on subsequent queues. (0 or 1)
                            • --kernelTime Time for each kernel execution
                            • --queuesCount How many queues to create
                            • --switchCount How many switches form each iteration
                            • --workgroupCount Count of workgroups
                            • --workgroupSize Size of the workgroup
                            		 ✔️ ❌
                            RoundTripSubmission enqueues kernel which updates system memory location and waits for it with a synchronizing API.
                              		✔️ ✔️
                              UniqueKernelSwitchLatency measures time to switch between unique kernels in a command list
                              • --kernelCount Count of unique kernel modules to execute, minimum 2
                              • --profiling Use profiling on the gpu to collect score (0 or 1)
                              • --workgroupCount Count of workgroups
                              • --workgroupSize Size of the workgroup
                              		 ✔️ ❌
                              UsmSharedFirstCpuAccess allocates a unified shared memory buffer and measures time to access it on CPU after creation.
                              • --initialPlacement Hint for initial placement of the resource passed to the driver (Any or Host or Device)
                              • --size Size of the buffer
                              		 ✔️ ✔️
                              UsmSharedFirstGpuAccess allocates a unified shared memory buffer and measures time to access it on GPU after creation.
                              • --initialPlacement Hint for initial placement of the resource passed to the driver (Any or Host or Device)
                              • --size Size of the buffer
                              		 ✔️ ✔️
                              WalkerCompletionLatency enqueues a kernel writing to system memory and measures time between the moment when update is visible on CPU and the moment when synchronizing call returns
                              • --inOrderQueue If set use IOQ, otherwise OOQ. Applicable only for OCL. (0 or 1)
                              • --useFence Use fence during submission and for further completion. (0 or 1)
                              		 ✔️ ✔️
                              WalkerSubmissionEvents enqueues an empty kernel with GPU-side profiling and checks delta between queue time and start time.
                                		✔️ ✔️
                                WriteLatency unblocks event on GPU, then waits for timestamp being written.
                                  		✔️ ❌
                                  ZMultiEngineSubmissionEvents creates one immediate command list for each engine from ccsMask and bcsMask, submits one command to all lists, and measures the submit delta. The delta is the time from the host API call to engine start. The test reports the delta for each engine and the worst value. If useHpCopyEngine is set, the test adds one high-priority copy engine list.
                                  • --bcsMask A bit mask for selecting copy engines, rightmost bit is the first copy engine reported by the driver
                                  • --ccsMask A bit mask for selecting compute engines, rightmost bit is CCS0
                                  • --useHpCopyEngine Adds one more list for the first copy engine reported by the driver, created with high priority (0 or 1)
                                  		 ✔️ ❌
                                  ZMultiQueueGraphSubmission replays a graph that forks from CCS0 to the engines from ccsMask and bcsMask, and measures the start latency of each branch and the CPU time of the replay call.
                                  • --bcsMask A bit mask for selecting copy engines the graph forks to, rightmost bit is the first copy engine reported by the driver. Engines that the device does not have are skipped
                                  • --ccsMask A bit mask for selecting compute engines the graph forks to, rightmost bit is CCS0. Engines that the device does not have are skipped
                                  • --useHpCopyEngine Adds one more branch for the first copy engine reported by the driver, on a list created with high priority (0 or 1)
                                  		 ✔️ ❌