Mutable Command List Extension#

API#

Mutable Command List#

A mutable command list is created by supplying a ze_mutable_command_list_exp_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_exp_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_exp_properties_t mutCmdListProps = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_LIST_EXP_PROPERTIES,      // stype
    nullptr,                                                    // pNext
    0,                                                          // mutableCommandListFlags
    0                                                           // mutableCommandFlags
};

ze_device_properties_t deviceProps = {
    ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES
};
deviceProps.pNext = &mutCmdListProps;

zeDeviceGetProperties(hDevice, &deviceProps);

// ...

// Create a mutable command list
ze_mutable_command_list_exp_desc_t mutCmdListDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_LIST_EXP_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_exp_desc_t cmdIdDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_COMMAND_ID_EXP_DESC       // stype
    nullptr,                                            // pNext
    0                                                   // flags
};

uint64_t commandId = 0;
zeCommandListGetNextCommandIdExp(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_exp_desc_t groupCountDesc = {
    ZE_STRUCTURE_TYPE_MUTABLE_GROUP_COUNT_EXP_DESC,         // stype
    nullptr,                                                // pNext
    commandId,                                              // commandId
    &groupCount                                             // pGroupCount
};

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

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

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

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

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

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

// Update the wait events for the launch kernel command
zeCommandListUpdateMutableCommandWaitEventsExp(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.

Mutable Command List Extension

Contents

Mutable Command List Extension#

API#

Mutable Command List#