Mutable Command List Extension#

Introduced in Level Zero version 1.20

API#

Mutable Command List#

A mutable command list is created by supplying a ze_mutable_command_list_ext_desc_t object via the pNext member of ze_command_list_desc_t.
Mutable command lists support mutation to identified commands after being closed with zeCommandListClose.
Implementation support for mutable commands may be discovered by providing a ze_mutable_command_list_ext_properties_t object in the pNext member of ze_device_properties_t in a call to zeDeviceGetProperties.

// Discover mutable command list properties
ze_mutable_command_list_ext_properties_t mutCmdListProps = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_LIST_EXT_PROPERTIES,      // stype
    nullptr,                                                    // pNext
    0,                                                          // mutableCommandListFlags
    0                                                           // mutableCommandFlags
};

ze_device_properties_t deviceProps = {
    ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES_1_2   // ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES deprecated since 1.17
};
deviceProps.pNext = &mutCmdListProps;

zeDeviceGetProperties(hDevice, &deviceProps);

// ...

// Create a mutable command list
ze_mutable_command_list_ext_desc_t mutCmdListDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_LIST_EXT_DESC,
    nullptr,
    0 // flags
};

ze_command_list_desc_t commandListDesc = {
    ZE_STRUCTURE_TYPE_COMMAND_LIST_DESC,
    &mutCmdListDesc,
    0,
    0 // flags
};

ze_command_list_handle_t hCommandList = nullptr;
zeCommandListCreate(hContext, hDevice, &commandListDesc, &hCommandList);

// [ ...create fence, signal event and wait event objects... ]
// This example assumes hFence, hSignalEvent and hWaitEvent have been created.

// Create kernel from module
ze_kernel_desc_t kernelDesc = {
    ZE_STRUCTURE_TYPE_KERNEL_DESC,
    0,
    "example"
};
ze_kernel_handle_t hKernel = nullptr;
ze_result_t result = zeKernelCreate(hModule, &kernelDesc, &hKernel);

// Set the kernel arguments
ze_group_count_t groupSize = {};
zeKernelSuggestGroupSize(hKernel, 1024, 1024, 1, &groupSize.groupCountX, &groupSize.groupCountY, &groupSize.groupCountZ);

int defaultValue = 0;
zeKernelSetArgumentValue(hKernel, 0, sizeof(int), &defaultValue);

// Get next command identifier
ze_mutable_command_id_ext_desc_t cmdIdDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_ID_EXT_DESC,      // stype
    nullptr,                                            // pNext
    0                                                   // flags
};

uint64_t commandId = 0;
zeCommandListGetNextCommandIdExt(hCommandList, &cmdIdDesc, &commandId);

// Encode command into command list
zeCommandListAppendLaunchKernel(hCommandList, hKernel, &groupSize, hSignalEvent, 1, &hWaitEvent);

// Close the command list
zeCommandListClose(hCommandList);

// ...

// Execute the command list
zeCommandQueueExecuteCommandLists(hCommandQueue, 1, &hCommandList, hFence);

// ...

The application may subsequently mutate specific commands, as follows:

// Prepare to modify group count
ze_group_count_t groupCount = {
    256,                                                    // groupCountX
    256,                                                    // groupCountY
    1                                                       // groupCountZ
};

ze_mutable_group_count_ext_desc_t groupCountDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_GROUP_COUNT_EXT_DESC,         // stype
    nullptr,                                                // pNext
    commandId,                                              // commandId
    &groupCount                                             // pGroupCount
};

// Prepare to modify Kernel Argument
int argValue = 1;

ze_mutable_kernel_argument_ext_desc_t krnlArgDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_KERNEL_ARGUMENT_EXT_DESC,     // stype
    &groupCountDesc,                                        // pNext
    commandId,                                              // commandId
    0,                                                      // argIndex
    sizeof(int),                                            // argSize
    &argValue                                               // pArgValue
};

// Prepare to update mutable commands
ze_mutable_commands_ext_desc_t desc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMANDS_EXT_DESC,            // stype
    &krnlArgDesc,                                           // pNext
    0                                                       // flags
};

// Synchronize command list execution
zeFenceHostSynchronize(hFence, UINT64_MAX);

// Update mutable commands
zeCommandListUpdateMutableCommandsExt(hCommandList, &desc);

// Update signal event for the launch kernel command
zeCommandListUpdateMutableCommandSignalEventExt(hCommandList, commandId, hNewLaunchKernelSignalEvent);

// Update the wait events for the launch kernel command
zeCommandListUpdateMutableCommandWaitEventsExt(hCommandList, commandId, 1, &hNewLaunchKernelWaitEvent);

// Close the command list
zeCommandListClose(hCommandList);

// ...

Note, the command list must be explicitly closed after updating mutable commands and events. This informs the implementation that the application has finished with updates and is ready to submit the command list. In preparation for kernel mutation user must provide all possible kernels for the command.

// define all possible kernels
ze_kernel_handle_t addKernel;
ze_kernel_handle_t mulKernel;

ze_kernel_handle_t kernels[] = {addKernel, mulKernel};

// when users want kernel mutation, they need to explicitly state this, as 0 does not include kernel instruction mutation by default
ze_mutable_command_ext_flags_t mutationFlags =
    ZE_MUTABLE_COMMAND_EXT_FLAG_KERNEL_ARGUMENTS |
    ZE_MUTABLE_COMMAND_EXT_FLAG_GROUP_COUNT |
    ZE_MUTABLE_COMMAND_EXT_FLAG_GROUP_SIZE |
    ZE_MUTABLE_COMMAND_EXT_FLAG_KERNEL_INSTRUCTION;

// Get next command identifier
ze_mutable_command_id_ext_desc_t cmdIdDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_ID_EXT_DESC,      // stype
    nullptr,                                            // pNext
    mutationFlags                                       // flags
};

// retrieve id for the append operation and provide all possible kernels for this command
uint64_t mutableKernelCommandId = 0;
zeCommandListGetNextCommandIdWithKernelsExt(hCommandList, &cmdIdDesc, &mutableKernelCommandId, 2, kernels);

// Encode command into command list
zeCommandListAppendLaunchKernel(hCommandList, addKernel, &groupSize, nullptr, 0, nullptr);

// Close the command list
zeCommandListClose(hCommandList);

Mutation of kernels must obey two rules: - kernel handle mutation function must be called as first for a given command id - kernel mutation invalidates all kernel arguments and dispatch parameters, these must be provided for the new kernel

// Update mutable kernel for the command, switch from addKernel to mulKernel
zeCommandListUpdateMutableCommandKernelsExt(hCommandList, 1, &mutableKernelCommandId, &mulKernel);

// modify group count
ze_group_count_t groupCount = {
    32,                                                     // groupCountX
    1,                                                      // groupCountY
    1                                                       // groupCountZ
};

ze_mutable_group_count_ext_desc_t groupCountDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_GROUP_COUNT_EXT_DESC,         // stype
    nullptr,                                                // pNext
    mutableKernelCommandId,                                 // commandId
    &groupCount                                             // pGroupCount
};

ze_mutable_group_size_ext_desc_t groupSizeDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_GROUP_SIZE_EXT_DESC,          // stype
    &groupCountDesc,                                        // pNext
    mutableKernelCommandId,                                 // commandId
    32,                                                     // groupSizeX
    1,                                                      // groupSizeY
    1,                                                      // groupSizeZ
};

// Prepare to modify Kernel Argument
int argValue = 1;
void *usmPointer;

ze_mutable_kernel_argument_ext_desc_t krnlArgMemoryDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_KERNEL_ARGUMENT_EXT_DESC,     // stype
    &groupSizeDesc,                                         // pNext
    mutableKernelCommandId,                                 // commandId
    0,                                                      // argIndex
    sizeof(void *),                                         // argSize
    &usmPointer                                             // pArgValue
};

ze_mutable_kernel_argument_ext_desc_t krnlArgScalarDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_KERNEL_ARGUMENT_EXT_DESC,     // stype
    &krnlArgMemoryDesc,                                     // pNext
    mutableKernelCommandId,                                 // commandId
    1,                                                      // argIndex
    sizeof(int),                                            // argSize
    &argValue                                               // pArgValue
};

// Prepare to update mutable commands
ze_mutable_commands_ext_desc_t desc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMANDS_EXT_DESC,            // stype
    &krnlArgScalarDesc,                                     // pNext
    0                                                       // flags
};

// Update mutable kernel arguments and dispatch parameters for the command
zeCommandListUpdateMutableCommandsExt(hCommandList, &desc);

// Close the command list
zeCommandListClose(hCommandList);

The command list must be explicitly closed after updating mutable commands.

Mutable Command List Extension

Contents

Mutable Command List Extension#

API#

Mutable Command List#