Mutable Command List Extension
Contents
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.