Add option to fill images on device

rjodinchr · rjodinchr · commit b283ed3d58ca · 2026-02-13T14:47:04.000+01:00
Create a second image (with write_only flag) using the same
VkDeviceMemory at runtime when it detects a fillimage on a read_only
image.
diff --git a/src/api.cpp b/src/api.cpp
@@ -5450,6 +5450,19 @@ cl_int CLVK_API_CALL clEnqueueFillImage(
     // Create image map command
     std::array<size_t, 3> reg = {region[0], region[1], region[2]};
 
+    if (config.fill_image_on_device()) {
+        std::array<size_t, 3> org = {origin[0], origin[1], origin[2]};
+        std::array<uint8_t, 16> fill_color_array;
+        std::memcpy(fill_color_array.data(), fill_color,
+                    sizeof(fill_color_array));
+
+        auto cmd_fill_on_device = new cvk_command_fill_image_on_device(
+            command_queue, image, fill_color_array, org, reg);
+        return command_queue->enqueue_command_with_deps(cmd_fill_on_device,
+                                                        num_events_in_wait_list,
+                                                        event_wait_list, event);
+    }
+
     void* map_ptr;
     _cl_event* evt_map;
     size_t image_row_pitch, image_slice_pitch;
diff --git a/src/config.def b/src/config.def
@@ -71,6 +71,8 @@ OPTION(bool, keep_memory_allocations_mapped, false)
 OPTION(std::string, device_extensions, "")
 OPTION(std::string, device_extensions_masked, "")
 
+OPTION(bool, fill_image_on_device, false)
+
 //
 // Logging
 //
diff --git a/src/device.cpp b/src/device.cpp
@@ -1292,10 +1292,12 @@ void cvk_device::select_work_group_size(
     cvk_kernel* kernel, const std::array<uint32_t, 3>& global_size,
     std::array<uint32_t, 3>& local_size) const {
 
-    auto required_work_group_size = kernel->required_work_group_size();
-    if (required_work_group_size[0] != 0) {
-        local_size = required_work_group_size;
-        return;
+    if (kernel != nullptr) {
+        auto required_work_group_size = kernel->required_work_group_size();
+        if (required_work_group_size[0] != 0) {
+            local_size = required_work_group_size;
+            return;
+        }
     }
     // Start at (1,1,1), which is always valid.
     local_size = {1, 1, 1};
diff --git a/src/memory.cpp b/src/memory.cpp
@@ -325,7 +325,21 @@ cvk_image* cvk_image::create(cvk_context* ctx, cl_mem_flags flags,
     return image.release();
 }
 
-bool cvk_image::init_vulkan_image() {
+cvk_image* cvk_image::create_write_enable_image_from(cvk_image* image) {
+    auto properties = image->properties();
+    auto image_write_enabled =
+        std::make_unique<cvk_image>(image->m_context, CL_MEM_WRITE_ONLY,
+                                    (const cl_image_desc*)&(image->m_desc),
+                                    (const cl_image_format*)&(image->m_format),
+                                    nullptr, std::move(properties));
+    if (!image_write_enabled->init(image->m_memory)) {
+        return nullptr;
+    }
+    return image_write_enabled.release();
+}
+
+bool cvk_image::init_vulkan_image(
+    std::shared_ptr<cvk_memory_allocation> memory) {
     // Translate image type and size
     VkImageType image_type;
     VkImageViewType view_type;
@@ -431,25 +445,29 @@ bool cvk_image::init_vulkan_image() {
         return false;
     }
 
-    CVK_ASSERT(m_desc.image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER);
-    // Select memory type
-    cvk_device::allocation_parameters params =
-        device->select_memory_for(m_image);
-    if (params.memory_type_index == VK_MAX_MEMORY_TYPES) {
-        cvk_error_fn("Could not get memory type!");
-        return false;
-    }
+    if (memory != VK_NULL_HANDLE) {
+        m_memory = memory;
+    } else {
+        CVK_ASSERT(m_desc.image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER);
+        // Select memory type
+        cvk_device::allocation_parameters params =
+            device->select_memory_for(m_image);
+        if (params.memory_type_index == VK_MAX_MEMORY_TYPES) {
+            cvk_error_fn("Could not get memory type!");
+            return false;
+        }
 
-    // Allocate memory
-    m_memory = std::make_unique<cvk_memory_allocation>(
-        vkdev, params.size, params.memory_type_index, params.memory_coherent,
-        device->keep_memory_allocations_mapped());
+        // Allocate memory
+        m_memory = std::make_unique<cvk_memory_allocation>(
+            vkdev, params.size, params.memory_type_index,
+            params.memory_coherent, device->keep_memory_allocations_mapped());
 
-    res = m_memory->allocate(device->uses_physical_addressing());
+        res = m_memory->allocate(device->uses_physical_addressing());
 
-    if (res != VK_SUCCESS) {
-        cvk_error_fn("Could not allocate memory!");
-        return false;
+        if (res != VK_SUCCESS) {
+            cvk_error_fn("Could not allocate memory!");
+            return false;
+        }
     }
 
     // Bind the image to memory
@@ -582,11 +600,11 @@ bool cvk_image::init_vulkan_texel_buffer() {
     return true;
 }
 
-bool cvk_image::init() {
+bool cvk_image::init(std::shared_ptr<cvk_memory_allocation> memory) {
     if (is_backed_by_buffer_view()) {
         return init_vulkan_texel_buffer();
     } else {
-        return init_vulkan_image();
+        return init_vulkan_image(memory);
     }
 }
 
diff --git a/src/memory.hpp b/src/memory.hpp
@@ -638,6 +638,7 @@ struct cvk_image : public cvk_mem {
                              const cl_image_desc* desc,
                              const cl_image_format* format, void* host_ptr,
                              std::vector<cl_mem_properties>&& properties);
+    static cvk_image* create_write_enable_image_from(cvk_image* image);
 
     bool is_backed_by_buffer_view() const {
         return type() == CL_MEM_OBJECT_IMAGE1D_BUFFER;
@@ -700,6 +701,7 @@ struct cvk_image : public cvk_mem {
     cvk_mem* buffer() const { return icd_downcast(m_desc.buffer); }
     cl_uint num_mip_levels() const { return m_desc.num_mip_levels; }
     cl_uint num_samples() const { return m_desc.num_samples; }
+    cl_image_info image_type() const { return m_desc.image_type; }
 
     bool has_same_format(const cvk_image* other) const {
         auto fmt = format();
@@ -828,9 +830,9 @@ struct cvk_image : public cvk_mem {
                               size_t* size_ret) const;
 
 private:
-    bool init_vulkan_image();
+    bool init_vulkan_image(std::shared_ptr<cvk_memory_allocation> memory);
     bool init_vulkan_texel_buffer();
-    bool init();
+    bool init(std::shared_ptr<cvk_memory_allocation> memory = nullptr);
 
     size_t num_channels() const {
         switch (m_format.image_channel_order) {
diff --git a/src/queue.cpp b/src/queue.cpp
@@ -1891,3 +1891,178 @@ cvk_command_image_init::build_batchable_inner(cvk_command_buffer& cmdbuf) {
 
     return CL_SUCCESS;
 }
+
+cl_int cvk_command_fill_image_on_device::build_batchable_inner(
+    cvk_command_buffer& cmdbuf) {
+    if (m_cmd_kernel != nullptr) {
+        return m_cmd_kernel->build_batchable_inner(cmdbuf);
+    }
+    char source[1024];
+    char color[512];
+    const char* kernel_name = "fill_image";
+    const char* type = get_image_type();
+    const char* access_qualifier = get_image_access_qualifier();
+    const char* coord = get_image_coord();
+    get_image_color(color);
+    sprintf(source,
+            "kernel void %s(write_only %s image) { write_image%s(image, "
+            "%s, %s); }",
+            kernel_name, type, access_qualifier, coord, color);
+
+    auto program = std::make_unique<cvk_program>(m_queue->context());
+    program->append_source(source, strlen(source));
+    const cl_device_id device = static_cast<cl_device_id>(m_queue->device());
+    auto err = program->build(build_operation::build, 1, &device, nullptr, 0,
+                              nullptr, nullptr, nullptr, nullptr);
+    if (err != CL_SUCCESS) {
+        cvk_error_fn("fail to build program (%u)", err);
+        return CL_OUT_OF_RESOURCES;
+    }
+
+    auto prog = program.release();
+    auto kernel = std::make_unique<cvk_kernel>(prog, kernel_name);
+    prog->release();
+    err = kernel->init();
+    if (err != CL_SUCCESS) {
+        cvk_error_fn("fail to init kernel (%u)", err);
+        return err;
+    }
+    err = kernel->set_arg(0, sizeof(m_mem), &m_mem);
+    if (err != CL_SUCCESS) {
+        cvk_error_fn("fail to set arg (%u)", err);
+        return err;
+    }
+
+    cl_uint work_dim = dimensions();
+    cvk_ndrange ndrange(work_dim, m_work_offset.data(), m_work_size.data(),
+                        nullptr);
+    auto kern = kernel.release();
+    m_cmd_kernel =
+        std::make_unique<cvk_command_kernel>(m_queue, kern, work_dim, ndrange);
+    kern->release();
+    return m_cmd_kernel->build_batchable_inner(cmdbuf);
+}
+
+cl_uint cvk_command_fill_image_on_device::dimensions() const {
+    switch (m_image->image_type()) {
+    default:
+        CVK_ASSERT(false);
+        return 0;
+    case CL_MEM_OBJECT_IMAGE1D:
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+        return 1;
+    case CL_MEM_OBJECT_IMAGE1D_ARRAY:
+    case CL_MEM_OBJECT_IMAGE2D:
+        return 2;
+    case CL_MEM_OBJECT_IMAGE2D_ARRAY:
+    case CL_MEM_OBJECT_IMAGE3D:
+        return 3;
+    }
+}
+
+const char*
+cvk_command_fill_image_on_device::get_image_access_qualifier() const {
+    switch (m_image->format().image_channel_data_type) {
+    default:
+        CVK_ASSERT(false);
+        return nullptr;
+    case CL_UNORM_SHORT_565:
+    case CL_UNORM_SHORT_555:
+    case CL_UNORM_INT_101010:
+    case CL_UNORM_INT_101010_2:
+    case CL_UNORM_INT8:
+    case CL_UNORM_INT16:
+        return "ui";
+    case CL_UNSIGNED_INT8:
+    case CL_UNSIGNED_INT16:
+    case CL_UNSIGNED_INT32:
+    case CL_SIGNED_INT8:
+    case CL_SIGNED_INT16:
+    case CL_SIGNED_INT32:
+    case CL_FLOAT:
+    case CL_HALF_FLOAT:
+        return "f";
+    case CL_SNORM_INT8:
+    case CL_SNORM_INT16:
+        return "i";
+    }
+}
+
+const char* cvk_command_fill_image_on_device::get_image_type() const {
+    switch (m_image->image_type()) {
+    default:
+        CVK_ASSERT(false);
+        return nullptr;
+    case CL_MEM_OBJECT_IMAGE1D:
+        return "image1d_t";
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+        return "image1d_buffer_t";
+    case CL_MEM_OBJECT_IMAGE1D_ARRAY:
+        return "image1d_array_t";
+    case CL_MEM_OBJECT_IMAGE2D:
+        return "image2d_t";
+    case CL_MEM_OBJECT_IMAGE2D_ARRAY:
+        return "image2d_array_t";
+    case CL_MEM_OBJECT_IMAGE3D:
+        return "image3d";
+    }
+}
+
+const char* cvk_command_fill_image_on_device::get_image_coord() const {
+    switch (m_image->image_type()) {
+    default:
+        CVK_ASSERT(false);
+        return nullptr;
+    case CL_MEM_OBJECT_IMAGE1D:
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+        return "get_global_id(0)";
+    case CL_MEM_OBJECT_IMAGE1D_ARRAY:
+    case CL_MEM_OBJECT_IMAGE2D:
+        return "(int2)(get_global_id(0), get_global_id(1))";
+    case CL_MEM_OBJECT_IMAGE2D_ARRAY:
+    case CL_MEM_OBJECT_IMAGE3D:
+        return "(int4)(get_global_id(0), get_global_id(1), "
+               "get_global_id(2), 0)";
+    }
+}
+
+void cvk_command_fill_image_on_device::get_image_color(char* color) const {
+    switch (m_image->format().image_channel_data_type) {
+    default:
+        CVK_ASSERT(false);
+        break;
+    case CL_UNORM_SHORT_565:
+    case CL_UNORM_SHORT_555:
+    case CL_UNORM_INT_101010:
+    case CL_UNORM_INT_101010_2:
+    case CL_UNORM_INT8:
+    case CL_UNORM_INT16: {
+        auto fill_color = static_cast<const uint32_t*>(
+            static_cast<const void*>(m_fill_color.data()));
+        sprintf(color, "(uint4)(%u, %u, %u, %u)", fill_color[0], fill_color[1],
+                fill_color[2], fill_color[3]);
+    } break;
+    case CL_UNSIGNED_INT8:
+    case CL_UNSIGNED_INT16:
+    case CL_UNSIGNED_INT32:
+    case CL_SIGNED_INT8:
+    case CL_SIGNED_INT16:
+    case CL_SIGNED_INT32:
+    case CL_FLOAT:
+    case CL_HALF_FLOAT: {
+        auto fill_color = static_cast<const uint32_t*>(
+            static_cast<const void*>(m_fill_color.data()));
+        sprintf(color,
+                "(float4)(as_float(0x%x), as_float(0x%x), as_float(0x%x), "
+                "as_float(0x%x))",
+                fill_color[0], fill_color[1], fill_color[2], fill_color[3]);
+    } break;
+    case CL_SNORM_INT8:
+    case CL_SNORM_INT16: {
+        auto fill_color = static_cast<const int32_t*>(
+            static_cast<const void*>(m_fill_color.data()));
+        sprintf(color, "(int4)(%i, %i, %i, %i)", fill_color[0], fill_color[1],
+                fill_color[2], fill_color[3]);
+    } break;
+    }
+}
diff --git a/src/queue.hpp b/src/queue.hpp
@@ -1115,3 +1115,59 @@ struct cvk_command_image_init final : public cvk_command_batchable {
 private:
     cvk_image_holder m_image;
 };
+
+struct cvk_command_fill_image_on_device final : public cvk_command_batchable {
+    cvk_command_fill_image_on_device(cvk_command_queue* queue, cl_mem image,
+                                     const std::array<uint8_t, 16> fill_color,
+                                     const std::array<size_t, 3> work_offset,
+                                     const std::array<size_t, 3> work_size)
+        : cvk_command_batchable(CL_COMMAND_FILL_IMAGE, queue),
+          m_image(static_cast<cvk_image*>(icd_downcast(image))),
+          m_fill_color(fill_color), m_work_offset(work_offset),
+          m_work_size(work_size) {
+        const std::array<uint32_t, 3> global_size = {(uint32_t)work_size[0],
+                                                     (uint32_t)work_size[1],
+                                                     (uint32_t)work_size[2]};
+        std::array<uint32_t, 3> local_size;
+        queue->device()->select_work_group_size(nullptr, global_size,
+                                                local_size);
+        m_local_size[0] = local_size[0];
+        m_local_size[1] = local_size[1];
+        m_local_size[2] = local_size[2];
+
+        if (!m_image->has_any_flag(CL_MEM_WRITE_ONLY | CL_MEM_READ_WRITE)) {
+            m_image_write_enabled.reset(
+                cvk_image::create_write_enable_image_from(m_image));
+            m_mem = m_image_write_enabled;
+        } else {
+            m_mem = image;
+        }
+    }
+
+    CHECK_RETURN cl_int
+    build_batchable_inner(cvk_command_buffer& cmdbuf) override final;
+
+    bool can_be_batched() const override final {
+        return cvk_command_batchable::can_be_batched();
+    }
+
+    const std::vector<cvk_mem*> memory_objects() const override {
+        return {m_image};
+    }
+
+private:
+    cl_uint dimensions() const;
+    const char* get_image_access_qualifier() const;
+    const char* get_image_type() const;
+    const char* get_image_coord() const;
+    void get_image_color(char* color) const;
+
+    cl_mem m_mem;
+    cvk_image_holder m_image;
+    cvk_image_holder m_image_write_enabled;
+    std::array<uint8_t, 16> m_fill_color;
+    std::array<size_t, 3> m_work_offset;
+    std::array<size_t, 3> m_work_size;
+    std::array<size_t, 3> m_local_size;
+    std::unique_ptr<cvk_command_kernel> m_cmd_kernel;
+};
