Level Zero: Question about host memory

[maleadt]

I have an issue with zeMemAllocHost-allocated memory incombination with a memset-like kernel. If I allocate the host memory, memset it on the device, copy it back using zeCommandListAppendMemoryCopy, everything is OK. If however instead of enqueuing a memory copy, I synchronize the command stream and use a regular CPU-based memcpy, I get the wrong results. Isn't host memory supposed to be readable from arbitrary CPU code? If it is, maybe the driver isn't fully synchronizing the command stream when calling zeCommandQueueSynchronize?

Here's a reproducer:

#include <level_zero/ze_api.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define CHECK_ERROR(x)                                                         \
  {                                                                            \
    if (x != ZE_RESULT_SUCCESS) {                                              \
      ze_result_t error = x;                                                   \
      const char *errorString;                                                 \
      zeDriverGetLastErrorDescription(NULL, &errorString);                     \
      fprintf(stderr, "Error: %d - %s\n", error, errorString);                 \
      abort();                                                                 \
    }                                                                          \
  }

char *read_binary_file(const char *filename, size_t *size) {
  FILE *file = fopen(filename, "rb");
  if (!file) {
    fprintf(stderr, "Failed to open file: %s\n", filename);
    return NULL;
  }

  fseek(file, 0, SEEK_END);
  *size = ftell(file);
  fseek(file, 0, SEEK_SET);

  char *buffer = (char *)malloc(*size);
  if (!buffer) {
    fclose(file);
    return NULL;
  }

  fread(buffer, 1, *size, file);
  fclose(file);
  return buffer;
}

int main() {
  // Initialize the driver
  CHECK_ERROR(zeInit(0));
  uint32_t driverCount = 0;
  CHECK_ERROR(zeDriverGet(&driverCount, NULL));
  ze_driver_handle_t driver;
  CHECK_ERROR(zeDriverGet(&driverCount, &driver));

  // Create a context
  ze_context_desc_t contextDesc = {ZE_STRUCTURE_TYPE_CONTEXT_DESC, NULL, 0};
  ze_context_handle_t context;
  CHECK_ERROR(zeContextCreate(driver, &contextDesc, &context));

  // Get the device
  uint32_t deviceCount = 0;
  CHECK_ERROR(zeDeviceGet(driver, &deviceCount, NULL));
  ze_device_handle_t device;
  CHECK_ERROR(zeDeviceGet(driver, &deviceCount, &device));

  // Create a command queue
  ze_command_queue_desc_t queueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC,
                                       NULL,
                                       0,
                                       0,
                                       0,
                                       ZE_COMMAND_QUEUE_MODE_DEFAULT,
                                       ZE_COMMAND_QUEUE_PRIORITY_NORMAL};
  ze_command_queue_handle_t queue;
  CHECK_ERROR(zeCommandQueueCreate(context, device, &queueDesc, &queue));

  // Allocate host memory
  const size_t size = sizeof(int);
  ze_host_mem_alloc_desc_t hostDesc = {ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC,
                                       NULL, 0};
  void *hostPtr;
  CHECK_ERROR(zeMemAllocHost(context, &hostDesc, size, 1, &hostPtr));

  // Create a module from SPIR-V
  size_t binary_size;
  char *binary = read_binary_file("kernel.spv", &binary_size);
  if (!binary) {
    fprintf(stderr, "Failed to read SPIR-V binary\n");
    return 1;
  }
  ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC,
                                 NULL,
                                 ZE_MODULE_FORMAT_IL_SPIRV,
                                 binary_size,
                                 binary,
                                 NULL,
                                 NULL};
  ze_module_handle_t module;
  ze_module_build_log_handle_t buildLog;
  CHECK_ERROR(zeModuleCreate(context, device, &moduleDesc, &module, &buildLog));

  // Prepare the kernel
  ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, NULL, 0,
                                 "_Z6memsetP5Int64S_"};
  ze_kernel_handle_t kernel;
  CHECK_ERROR(zeKernelCreate(module, &kernelDesc, &kernel));
  int value = 42;
  CHECK_ERROR(zeKernelSetArgumentValue(kernel, 0, sizeof(void *), &hostPtr));
  CHECK_ERROR(zeKernelSetArgumentValue(kernel, 1, sizeof(int), &value));
  CHECK_ERROR(zeKernelSetGroupSize(kernel, 1, 1, 1));

  // Launch the kernel using a command list
  ze_command_list_desc_t cmdListDesc = {ZE_STRUCTURE_TYPE_COMMAND_LIST_DESC,
                                        NULL, 0, 0};
  ze_command_list_handle_t cmdList;
  CHECK_ERROR(zeCommandListCreate(context, device, &cmdListDesc, &cmdList));
  ze_group_count_t launchArgs = {1, 1, 1};
  CHECK_ERROR(zeCommandListAppendLaunchKernel(cmdList, kernel, &launchArgs,
                                              NULL, 0, NULL));
  CHECK_ERROR(zeCommandListClose(cmdList));
  CHECK_ERROR(zeCommandQueueExecuteCommandLists(queue, 1, &cmdList, NULL));

  // Synchronize
  CHECK_ERROR(zeCommandQueueSynchronize(queue, UINT64_MAX));

  // Print result
  printf("Result: %d\n", *(int *)hostPtr);

  // Cleanup
  zeMemFree(context, hostPtr);
  zeCommandListDestroy(cmdList);
  zeKernelDestroy(kernel);
  zeModuleDestroy(module);
  zeCommandQueueDestroy(queue);
  zeContextDestroy(context);
  free(binary);

  return 0;
}

This reads kernel.spv, attached here: kernel.zip

Originally reported in: JuliaGPU/oneAPI.jl#458 (comment)