From e9f01cf26887b057ad912b29ae1077eaaa10c34a Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Sun, 7 Sep 2025 01:04:15 +0200 Subject: [PATCH] adding device selection, buffer creation and pipeline creation to Metal backend --- .github/workflows/metal-test-macos.yml | 86 +++++++++++++ Examples/Metal/main.c | 63 ++++++++++ Examples/Metal/xmake.lua | 4 + Examples/xmake.lua | 5 +- Includes/PulseProfile.h | 17 ++- Sources/Backends/Metal/Metal.h | 17 ++- Sources/Backends/Metal/Metal.m | 22 +++- Sources/Backends/Metal/MetalBuffer.h | 29 +++++ Sources/Backends/Metal/MetalBuffer.m | 73 +++++++++++ Sources/Backends/Metal/MetalCommandList.h | 29 +++++ Sources/Backends/Metal/MetalCommandList.m | 29 +++++ Sources/Backends/Metal/MetalComputePass.h | 27 ++++ Sources/Backends/Metal/MetalComputePass.m | 49 ++++++++ Sources/Backends/Metal/MetalComputePipeline.h | 25 ++++ Sources/Backends/Metal/MetalComputePipeline.m | 109 +++++++++++++++++ Sources/Backends/Metal/MetalDevice.h | 26 ++++ Sources/Backends/Metal/MetalDevice.m | 115 ++++++++++++++++++ Sources/Backends/Metal/MetalFence.h | 27 ++++ Sources/Backends/Metal/MetalFence.m | 31 +++++ Sources/Backends/Metal/MetalImage.h | 27 ++++ Sources/Backends/Metal/MetalImage.m | 28 +++++ Sources/Backends/Vulkan/VulkanDevice.c | 8 +- Sources/Backends/Vulkan/VulkanInstance.c | 4 +- Sources/Backends/Vulkan/VulkanLoader.c | 2 +- Sources/PulseDefs.h | 2 + Sources/PulseInternal.c | 2 +- xmake.lua | 8 +- 27 files changed, 848 insertions(+), 16 deletions(-) create mode 100644 .github/workflows/metal-test-macos.yml create mode 100644 Examples/Metal/main.c create mode 100644 Examples/Metal/xmake.lua create mode 100644 Sources/Backends/Metal/MetalBuffer.h create mode 100644 Sources/Backends/Metal/MetalBuffer.m create mode 100644 Sources/Backends/Metal/MetalCommandList.h create mode 100644 Sources/Backends/Metal/MetalCommandList.m create mode 100644 Sources/Backends/Metal/MetalComputePass.h create mode 100644 Sources/Backends/Metal/MetalComputePass.m create mode 100644 Sources/Backends/Metal/MetalComputePipeline.h create mode 100644 Sources/Backends/Metal/MetalComputePipeline.m create mode 100644 Sources/Backends/Metal/MetalDevice.h create mode 100644 Sources/Backends/Metal/MetalDevice.m create mode 100644 Sources/Backends/Metal/MetalFence.h create mode 100644 Sources/Backends/Metal/MetalFence.m create mode 100644 Sources/Backends/Metal/MetalImage.h create mode 100644 Sources/Backends/Metal/MetalImage.m diff --git a/.github/workflows/metal-test-macos.yml b/.github/workflows/metal-test-macos.yml new file mode 100644 index 0000000..1201333 --- /dev/null +++ b/.github/workflows/metal-test-macos.yml @@ -0,0 +1,86 @@ +name: Metal + +on: + pull_request: + push: + paths-ignore: + - '.github/workflows/*.yml' + - '!.github/workflows/metal-test-macos.yml' + - '.gitignore' + - 'LICENSE' + - 'CHANGELOG.md' + - 'README.md' + +jobs: + build: + strategy: + fail-fast: false + matrix: + os: [macOS-latest] + arch: [x86_64] + confs: + - { mode: debug, archive: yes } + + runs-on: ${{ matrix.os }} + if: ${{ !contains(github.event.head_commit.message, 'ci skip') }} + + steps: + - name: Get current date as package key + id: cache_key + run: echo "key=$(date +'%W')" >> $GITHUB_OUTPUT + + - name: Checkout repository + uses: actions/checkout@v4 + + # Force xmake to a specific folder (for cache) + - name: Set xmake env + run: echo "XMAKE_GLOBALDIR=${{ runner.workspace }}/xmake-global" >> $GITHUB_ENV + + # Install xmake + - name: Setup xmake + uses: xmake-io/github-action-setup-xmake@v1 + with: + xmake-version: branch@dev + actions-cache-folder: .xmake-cache-W${{ steps.cache_key.outputs.key }} + + # Update xmake repository (in order to have the file that will be cached) + - name: Update xmake repository + run: xmake repo --update + + # Fetch xmake dephash + - name: Retrieve dependencies hash + id: dep_hash + run: echo "hash=$(xmake l utils.ci.packageskey)" >> $GITHUB_OUTPUT + + # Cache xmake dependencies + - name: Restore cached xmake dependencies + id: restore-depcache + uses: actions/cache/restore@v4 + with: + path: ${{ env.XMAKE_GLOBALDIR }}/.xmake/packages + key: MacOS-${{ matrix.arch }}-${{ matrix.confs.mode }}${{ matrix.confs.cache_key }}-${{ steps.dep_hash.outputs.hash }}-W${{ steps.cache_key.outputs.key }} + + # Setup compilation mode and install project dependencies + - name: Configure xmake and install dependencies + run: xmake config --examples=y --arch=${{ matrix.arch }} --mode=${{ matrix.confs.mode }} ${{ matrix.confs.config }} --ccache=n --yes + + # Save dependencies + - name: Save cached xmake dependencies + if: ${{ !steps.restore-depcache.outputs.cache-hit }} + uses: actions/cache/save@v4 + with: + path: ${{ env.XMAKE_GLOBALDIR }}/.xmake/packages + key: ${{ steps.restore-depcache.outputs.cache-primary-key }} + + # Cache assets downloading + - name: Restore cached assets + id: restore-assets + uses: actions/cache/restore@v4 + with: + path: assets + key: assets-${{ hashFiles('assets/examples_version.txt', 'assets/unittests_version.txt') }} + + - name: Test Metal + run: | + xmake build --yes MetalExample + xmake run --yes MetalExample diff --git a/Examples/Metal/main.c b/Examples/Metal/main.c new file mode 100644 index 0000000..71522e9 --- /dev/null +++ b/Examples/Metal/main.c @@ -0,0 +1,63 @@ +#include + +#include +#include +#include + +void DebugCallBack(PulseDebugMessageSeverity severity, const char* message) +{ + if(severity == PULSE_DEBUG_MESSAGE_SEVERITY_ERROR) + { + fprintf(stderr, "Pulse Error: %s\n", message); + exit(1); + } + else if(severity == PULSE_DEBUG_MESSAGE_SEVERITY_WARNING) + fprintf(stderr, "Pulse Warning: %s\n", message); + else + printf("Pulse: %s\n", message); +} + +#define BUFFER_SIZE (256 * sizeof(uint32_t)) + +const char* msl_source = +"#include \n \ +using namespace metal;\ +\ +kernel void Main(device int* ssbo [[ buffer(0) ]], uint3 grid [[ thread_position_in_grid ]]) \ +{ \ + uint idx = grid.x * grid.y; \ + ssbo[idx] = static_cast(idx); \ +}"; + +int main(void) +{ + PulseBackend backend = PulseLoadBackend(PULSE_BACKEND_METAL, PULSE_SHADER_FORMAT_MSL_BIT, PULSE_HIGH_DEBUG); + PulseSetDebugCallback(backend, DebugCallBack); + PulseDevice device = PulseCreateDevice(backend, NULL, 0); + + PulseBufferCreateInfo buffer_create_info = { 0 }; + buffer_create_info.size = BUFFER_SIZE; + buffer_create_info.usage = PULSE_BUFFER_USAGE_STORAGE_READ | PULSE_BUFFER_USAGE_STORAGE_WRITE | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD; + PulseBuffer buffer = PulseCreateBuffer(device, &buffer_create_info); + + // GPU computations + { + PulseComputePipelineCreateInfo info = { 0 }; + info.code_size = strlen(msl_source); + info.code = (const uint8_t*)msl_source; + info.entrypoint = "Main"; + info.format = PULSE_SHADER_FORMAT_MSL_BIT; + info.num_readwrite_storage_buffers = 1; + PulseComputePipeline pipeline = PulseCreateComputePipeline(device, &info); + + PulseDestroyComputePipeline(device, pipeline); + } + + PulseDestroyBuffer(device, buffer); + + PulseDestroyDevice(device); + PulseUnloadBackend(backend); + puts("Successfully executed Pulse example using Metal !"); + return 0; +} + diff --git a/Examples/Metal/xmake.lua b/Examples/Metal/xmake.lua new file mode 100644 index 0000000..9d17a5e --- /dev/null +++ b/Examples/Metal/xmake.lua @@ -0,0 +1,4 @@ +target("MetalExample") + add_deps("pulse_gpu") + add_files("*.c") +target_end() diff --git a/Examples/xmake.lua b/Examples/xmake.lua index 4c58db6..39ee0db 100644 --- a/Examples/xmake.lua +++ b/Examples/xmake.lua @@ -2,9 +2,12 @@ option("examples", { description = "Build the examples", default = false }) if has_config("examples") then set_group("Examples") - if not is_plat("wasm") and has_config("vulkan") then + if has_config("vulkan") then includes("Vulkan/xmake.lua") end + if has_config("metal") then + includes("Metal/xmake.lua") + end if has_config("webgpu") then includes("WebGPU/xmake.lua") end diff --git a/Includes/PulseProfile.h b/Includes/PulseProfile.h index a2fe0d7..c8396b1 100644 --- a/Includes/PulseProfile.h +++ b/Includes/PulseProfile.h @@ -48,8 +48,23 @@ extern "C" { #define PULSE_PLAT_LINUX #define PULSE_PLAT_POSIX #elif defined(__APPLE__) && defined(__MACH__) - #define PULSE_PLAT_MACOS + #define PULSE_PLAT_APPLE #define PULSE_PLAT_POSIX + #include + #if TARGET_IPHONE_SIMULATOR + // iOS, tvOS, or watchOS Simulator + #define PULSE_PLAT_IOS + #elif TARGET_OS_MACCATALYST + // Mac's Catalyst (ports iOS API into Mac, like UIKit). + #define PULSE_PLAT_IOS + #elif TARGET_OS_IPHONE + // iOS, tvOS, or watchOS device + #define PULSE_PLAT_IOS + #elif TARGET_OS_MAC + #define PULSE_PLAT_MACOS + #else + #error "Unknown Apple platform" + #endif #elif defined(unix) || defined(__unix__) || defined(__unix) #define PULSE_PLAT_UNIX #define PULSE_PLAT_POSIX diff --git a/Sources/Backends/Metal/Metal.h b/Sources/Backends/Metal/Metal.h index 8a2a5e6..3b27cc2 100644 --- a/Sources/Backends/Metal/Metal.h +++ b/Sources/Backends/Metal/Metal.h @@ -9,7 +9,22 @@ #ifndef PULSE_METAL_H_ #define PULSE_METAL_H_ -PulseBackendFlags MetalCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used); // Return PULSE_BACKEND_METAL in case of success and PULSE_BACKEND_INVALID otherwise +#define METAL_RETRIEVE_DRIVER_DATA_AS(handle, cast) ((cast)handle->driver_data) + +#define CHECK_METAL_RETVAL(backend, handle, error, retval) \ + do { \ + if(!(handle)) \ + { \ + if(backend != PULSE_NULL_HANDLE && PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend)) \ + PulseLogError(backend, "(Metal) call to a Metal function failed"); \ + PulseSetInternalError(error); \ + return retval; \ + } \ + } while(0) \ + +#define CHECK_METAL(backend, handle, error) CHECK_METAL_RETVAL(backend, handle, error, ) + +PulseBackendFlags MetalCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used); // Returns corresponding PULSE_BACKEND enum in case of success and PULSE_BACKEND_INVALID otherwise #endif // PULSE_METAL_H_ diff --git a/Sources/Backends/Metal/Metal.m b/Sources/Backends/Metal/Metal.m index fa7ec9a..e8287b5 100644 --- a/Sources/Backends/Metal/Metal.m +++ b/Sources/Backends/Metal/Metal.m @@ -6,20 +6,32 @@ #include "../../PulseInternal.h" #include "Metal.h" +#include "MetalDevice.h" PulseBackendFlags MetalCheckSupport(PulseBackendFlags candidates, PulseShaderFormatsFlags shader_formats_used) { if(candidates != PULSE_BACKEND_ANY && (candidates & PULSE_BACKEND_METAL) == 0) return PULSE_BACKEND_INVALID; - if((shader_formats_used & PULSE_SHADER_FORMAT_MSL_BIT) == 0 && (shader_formats_used & PULSE_SHADER_FORMAT_METALLIB_BIT) == 0) + if((shader_formats_used & (PULSE_SHADER_FORMAT_MSL_BIT | PULSE_SHADER_FORMAT_METALLIB_BIT)) == 0) return PULSE_BACKEND_INVALID; - return PULSE_BACKEND_INVALID; // Not supported yet + return PULSE_BACKEND_METAL; +} + +bool MetalLoadBackend(PulseBackend backend, PulseDebugLevel debug_level) +{ + PULSE_UNUSED(backend); + PULSE_UNUSED(debug_level); + return true; +} + +void MetalUnloadBackend(PulseBackend backend) +{ } PulseBackendHandler MetalDriver = { - .PFN_LoadBackend = PULSE_NULLPTR, - .PFN_UnloadBackend = PULSE_NULLPTR, - .PFN_CreateDevice = PULSE_NULLPTR, + .PFN_LoadBackend = MetalLoadBackend, + .PFN_UnloadBackend = MetalUnloadBackend, + .PFN_CreateDevice = MetalCreateDevice, .backend = PULSE_BACKEND_METAL, .supported_shader_formats = PULSE_SHADER_FORMAT_MSL_BIT | PULSE_SHADER_FORMAT_METALLIB_BIT, .driver_data = PULSE_NULLPTR diff --git a/Sources/Backends/Metal/MetalBuffer.h b/Sources/Backends/Metal/MetalBuffer.h new file mode 100644 index 0000000..554c6a8 --- /dev/null +++ b/Sources/Backends/Metal/MetalBuffer.h @@ -0,0 +1,29 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_BUFFER_H_ +#define PULSE_METAL_BUFFER_H_ + +#include +#include "Metal.h" + +typedef struct MetalBuffer +{ + id buffer; +} MetalBuffer; + +PulseBuffer MetalCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos); +bool MetalMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data); +void MetalUnmapBuffer(PulseBuffer buffer); +bool MetalCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst); +bool MetalCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst); +void MetalDestroyBuffer(PulseDevice device, PulseBuffer buffer); + +#endif // PULSE_METAL_BUFFER_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalBuffer.m b/Sources/Backends/Metal/MetalBuffer.m new file mode 100644 index 0000000..679acad --- /dev/null +++ b/Sources/Backends/Metal/MetalBuffer.m @@ -0,0 +1,73 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalBuffer.h" +#include "MetalDevice.h" +#include "MetalCommandList.h" + +PulseBuffer MetalCreateBuffer(PulseDevice device, const PulseBufferCreateInfo* create_infos) +{ + @autoreleasepool + { + MetalDevice* metal_device = (MetalDevice*)METAL_RETRIEVE_DRIVER_DATA_AS(device, MetalDevice*); + + PulseBuffer buffer = (PulseBuffer)calloc(1, sizeof(PulseBufferHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(buffer, PULSE_NULL_HANDLE); + + MetalBuffer* metal_buffer = (MetalBuffer*)calloc(1, sizeof(MetalBuffer)); + PULSE_CHECK_ALLOCATION_RETVAL(metal_buffer, PULSE_NULL_HANDLE); + + uint32_t size = PULSE_ALIGN_UP(create_infos->size, 4); + + buffer->device = device; + buffer->driver_data = metal_buffer; + buffer->size = size; + buffer->usage = create_infos->usage; + + MTLResourceOptions options; + if((create_infos->usage & (PULSE_BUFFER_USAGE_TRANSFER_UPLOAD | PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD)) == 0) // Is storage only + options = MTLResourceStorageModePrivate; + else if((create_infos->usage & PULSE_BUFFER_USAGE_TRANSFER_DOWNLOAD) == 0) + options = MTLResourceCPUCacheModeWriteCombined; + else + options = MTLResourceCPUCacheModeDefaultCache; + + metal_buffer->buffer = [metal_device->device newBufferWithLength:size options:options]; + CHECK_METAL_RETVAL(device->backend, metal_buffer->buffer, PULSE_ERROR_INITIALIZATION_FAILED, PULSE_NULL_HANDLE); + + return buffer; + } +} + +bool MetalMapBuffer(PulseBuffer buffer, PulseMapMode mode, void** data) +{ + return true; +} + +void MetalUnmapBuffer(PulseBuffer buffer) +{ +} + +bool MetalCopyBufferToBuffer(PulseCommandList cmd, const PulseBufferRegion* src, const PulseBufferRegion* dst) +{ + return true; +} + +bool MetalCopyBufferToImage(PulseCommandList cmd, const PulseBufferRegion* src, const PulseImageRegion* dst) +{ + return true; +} + +void MetalDestroyBuffer(PulseDevice device, PulseBuffer buffer) +{ + MetalBuffer* metal_buffer = METAL_RETRIEVE_DRIVER_DATA_AS(buffer, MetalBuffer*); + metal_buffer->buffer = nil; + free(metal_buffer); + free(buffer); +} diff --git a/Sources/Backends/Metal/MetalCommandList.h b/Sources/Backends/Metal/MetalCommandList.h new file mode 100644 index 0000000..bfc48ac --- /dev/null +++ b/Sources/Backends/Metal/MetalCommandList.h @@ -0,0 +1,29 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_COMMAND_LIST_H_ +#define PULSE_METAL_COMMAND_LIST_H_ + +#include +#include + +#include "Metal.h" +#include "MetalFence.h" + +typedef struct MetalCommandList +{ + int dummy; +} MetalCommandList; + +PulseCommandList MetalRequestCommandList(PulseDevice device, PulseCommandListUsage usage); +bool MetalSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence); +void MetalReleaseCommandList(PulseDevice device, PulseCommandList cmd); + +#endif // PULSE_METAL_COMMAND_LIST_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalCommandList.m b/Sources/Backends/Metal/MetalCommandList.m new file mode 100644 index 0000000..a5086b3 --- /dev/null +++ b/Sources/Backends/Metal/MetalCommandList.m @@ -0,0 +1,29 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalFence.h" +#include "MetalDevice.h" +#include "MetalCommandList.h" +#include "MetalComputePass.h" +#include "MetalComputePipeline.h" +#include "MetalBuffer.h" + +PulseCommandList MetalRequestCommandList(PulseDevice device, PulseCommandListUsage usage) +{ + PULSE_CHECK_HANDLE_RETVAL(device, PULSE_NULL_HANDLE); + PulseCommandList cmd = (PulseCommandList)calloc(1, sizeof(PulseCommandListHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(cmd, PULSE_NULL_HANDLE); + return cmd; +} + +bool MetalSubmitCommandList(PulseDevice device, PulseCommandList cmd, PulseFence fence) +{ +} + +void MetalReleaseCommandList(PulseDevice device, PulseCommandList cmd) +{ +} diff --git a/Sources/Backends/Metal/MetalComputePass.h b/Sources/Backends/Metal/MetalComputePass.h new file mode 100644 index 0000000..cbc77a0 --- /dev/null +++ b/Sources/Backends/Metal/MetalComputePass.h @@ -0,0 +1,27 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_COMPUTE_PASS_H_ +#define PULSE_METAL_COMPUTE_PASS_H_ + +#include "Metal.h" + +PulseComputePass MetalCreateComputePass(PulseDevice device, PulseCommandList cmd); +void MetalDestroyComputePass(PulseDevice device, PulseComputePass pass); + +PulseComputePass MetalBeginComputePass(PulseCommandList cmd); +void MetalEndComputePass(PulseComputePass pass); +void MetalBindStorageBuffers(PulseComputePass pass, const PulseBuffer* buffers, uint32_t num_buffers); +void MetalBindUniformData(PulseComputePass pass, uint32_t slot, const void* data, uint32_t data_size); +void MetalBindStorageImages(PulseComputePass pass, const PulseImage* images, uint32_t num_images); +void MetalBindComputePipeline(PulseComputePass pass, PulseComputePipeline pipeline); +void MetalDispatchComputations(PulseComputePass pass, uint32_t groupcount_x, uint32_t groupcount_y, uint32_t groupcount_z); + +#endif // PULSE_METAL_COMPUTE_PASS_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalComputePass.m b/Sources/Backends/Metal/MetalComputePass.m new file mode 100644 index 0000000..25b3aae --- /dev/null +++ b/Sources/Backends/Metal/MetalComputePass.m @@ -0,0 +1,49 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalComputePass.h" +#include "MetalCommandList.h" + +PulseComputePass MetalCreateComputePass(PulseDevice device, PulseCommandList cmd) +{ + PULSE_UNUSED(device); + PulseComputePass pass = (PulseComputePass)calloc(1, sizeof(PulseComputePassHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(pass, PULSE_NULL_HANDLE); + return pass; +} + +void MetalDestroyComputePass(PulseDevice device, PulseComputePass pass) +{ +} + +PulseComputePass MetalBeginComputePass(PulseCommandList cmd) +{ +} + +void MetalEndComputePass(PulseComputePass pass) +{ +} + +void MetalBindStorageBuffers(PulseComputePass pass, const PulseBuffer* buffers, uint32_t num_buffers) +{ +} + +void MetalBindUniformData(PulseComputePass pass, uint32_t slot, const void* data, uint32_t data_size) +{ +} + +void MetalBindStorageImages(PulseComputePass pass, const PulseImage* images, uint32_t num_images) +{ +} + +void MetalBindComputePipeline(PulseComputePass pass, PulseComputePipeline pipeline) +{ +} + +void MetalDispatchComputations(PulseComputePass pass, uint32_t groupcount_x, uint32_t groupcount_y, uint32_t groupcount_z) +{ +} diff --git a/Sources/Backends/Metal/MetalComputePipeline.h b/Sources/Backends/Metal/MetalComputePipeline.h new file mode 100644 index 0000000..b9cf960 --- /dev/null +++ b/Sources/Backends/Metal/MetalComputePipeline.h @@ -0,0 +1,25 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_COMPUTE_PIPELINE_H_ +#define PULSE_METAL_COMPUTE_PIPELINE_H_ + +#include +#include "Metal.h" + +typedef struct MetalComputePipeline +{ + id pipeline; +} MetalComputePipeline; + +PulseComputePipeline MetalCreateComputePipeline(PulseDevice device, const PulseComputePipelineCreateInfo* info); +void MetalDestroyComputePipeline(PulseDevice device, PulseComputePipeline pipeline); + +#endif // PULSE_METAL_COMPUTE_PIPELINE_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalComputePipeline.m b/Sources/Backends/Metal/MetalComputePipeline.m new file mode 100644 index 0000000..58e39f6 --- /dev/null +++ b/Sources/Backends/Metal/MetalComputePipeline.m @@ -0,0 +1,109 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalDevice.h" +#include "MetalComputePipeline.h" + +typedef struct MetalLibraryFunction +{ + id library; + id function; +} MetalLibraryFunction; + +static bool MetalIsValidMetalLibrary(const uint8_t* code, uint32_t code_size) +{ + // Metal libraries have a 4 byte header containing `MTLB`. + if(code_size < 4 || code == PULSE_NULLPTR) + return false; + return memcmp(code, "MTLB", 4) == 0; +} + +static MetalLibraryFunction MetalCompileShader(PulseDevice device, const PulseComputePipelineCreateInfo* info) +{ + MetalDevice* metal_device = (MetalDevice*)METAL_RETRIEVE_DRIVER_DATA_AS(device, MetalDevice*); + + MetalLibraryFunction library_function = { nil, nil }; + id library; + NSError* error; + + const char* entrypoint = (info->entrypoint == PULSE_NULLPTR) ? "main0" : info->entrypoint; + + if(info->format == PULSE_SHADER_FORMAT_MSL_BIT) + { + NSString* code_string = [[NSString alloc] initWithBytes:info->code length:info->code_size encoding:NSUTF8StringEncoding]; + library = [metal_device->device newLibraryWithSource:code_string options:nil error:&error]; + } + else if(info->format == PULSE_SHADER_FORMAT_METALLIB_BIT) + { + if(!MetalIsValidMetalLibrary(info->code, info->code_size)) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(device->backend)) + PulseLogError(device->backend, "(Metal) provided shader code is not a valid Metal library"); + return library_function; + } + dispatch_data_t data = dispatch_data_create(info->code, info->code_size, dispatch_get_global_queue(0, 0), DISPATCH_DATA_DESTRUCTOR_DEFAULT); + library = [metal_device->device newLibraryWithData:data error:&error]; + } + + if(library == nil) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(device->backend)) + PulseLogErrorFmt(device->backend, "(Metal) creating MTLLibrary failed due to %s", [[error description] cStringUsingEncoding:[NSString defaultCStringEncoding]]); + return library_function; + } + else if(error != nil && PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(device->backend)) + PulseLogWarningFmt(device->backend, "(Metal) creating MTLLibrary failed due to %s", [[error description] cStringUsingEncoding:[NSString defaultCStringEncoding]]); + + id function = [library newFunctionWithName:@(entrypoint)]; + if(function == nil) + { + PulseLogError(device->backend, "(Metal) creating MTLLibrary failed"); + return library_function; + } + + library_function.library = library; + library_function.function = function; + return library_function; +} + +PulseComputePipeline MetalCreateComputePipeline(PulseDevice device, const PulseComputePipelineCreateInfo* info) +{ + @autoreleasepool + { + PulseComputePipelineHandler* pipeline = (PulseComputePipelineHandler*)calloc(1, sizeof(PulseComputePipelineHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(pipeline, PULSE_NULL_HANDLE); + + MetalComputePipeline* metal_pipeline = (MetalComputePipeline*)calloc(1, sizeof(MetalComputePipeline)); + PULSE_CHECK_ALLOCATION_RETVAL(metal_pipeline, PULSE_NULL_HANDLE); + + MetalLibraryFunction library_function = MetalCompileShader(device, info); + if(library_function.library == nil || library_function.function == nil) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(device->backend)) + PulseLogError(device->backend, "(Metal) failed to compile shader"); + PulseSetInternalError(PULSE_ERROR_INITIALIZATION_FAILED); + return PULSE_NULL_HANDLE; + } + + pipeline->driver_data = metal_pipeline; + + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend)) + PulseLogInfoFmt(device->backend, "(Metal) created new compute pipeline %p", pipeline); + return pipeline; + } +} + +void MetalDestroyComputePipeline(PulseDevice device, PulseComputePipeline pipeline) +{ + MetalComputePipeline* metal_pipeline = METAL_RETRIEVE_DRIVER_DATA_AS(pipeline, MetalComputePipeline*); + metal_pipeline->pipeline = nil; + free(metal_pipeline); + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend)) + PulseLogInfoFmt(device->backend, "(Metal) destroyed compute pipeline %p", pipeline); + free(pipeline); +} diff --git a/Sources/Backends/Metal/MetalDevice.h b/Sources/Backends/Metal/MetalDevice.h new file mode 100644 index 0000000..874d066 --- /dev/null +++ b/Sources/Backends/Metal/MetalDevice.h @@ -0,0 +1,26 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_DEVICE_H_ +#define PULSE_METAL_DEVICE_H_ + +#include +#include "Metal.h" + +typedef struct MetalDevice +{ + id device; + id queue; +} MetalDevice; + +PulseDevice MetalCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count); +void MetalDestroyDevice(PulseDevice device); + +#endif // PULSE_METAL_DEVICE_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalDevice.m b/Sources/Backends/Metal/MetalDevice.m new file mode 100644 index 0000000..e1864f3 --- /dev/null +++ b/Sources/Backends/Metal/MetalDevice.m @@ -0,0 +1,115 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalComputePipeline.h" +#include "MetalCommandList.h" +#include "MetalDevice.h" +#include "MetalFence.h" +#include "MetalBuffer.h" +#include "MetalImage.h" +#include "MetalComputePass.h" + +static uint64_t MetalScoreDevice(id device) +{ + uint64_t score = 0; + + if(!device.lowPower) + score += 1000; + + // A GPU with dedicated memory is typically a dedicated one + if(!device.hasUnifiedMemory) + score += 10000; + + score += device.maxThreadsPerThreadgroup.width; + score += device.maxThreadsPerThreadgroup.height; + score += device.maxThreadsPerThreadgroup.depth; + score += device.maxThreadgroupMemoryLength; + + return score; +} + +static bool MetalIsDeviceForbidden(id device, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count) +{ + if(!device) + return true; + for(uint32_t i = 0; i < forbiden_devices_count; i++) + { + if(device.registryID == METAL_RETRIEVE_DRIVER_DATA_AS(forbiden_devices[i], MetalDevice*)->device.registryID) + return true; + } + return false; +} + +PulseDevice MetalCreateDevice(PulseBackend backend, PulseDevice* forbiden_devices, uint32_t forbiden_devices_count) +{ + @autoreleasepool + { + PULSE_UNUSED(forbiden_devices); + PULSE_UNUSED(forbiden_devices_count); + + PULSE_CHECK_HANDLE_RETVAL(backend, PULSE_NULLPTR); + + PulseDevice pulse_device = (PulseDeviceHandler*)calloc(1, sizeof(PulseDeviceHandler)); + PULSE_CHECK_ALLOCATION_RETVAL(pulse_device, PULSE_NULL_HANDLE); + + MetalDevice* metal_device = (MetalDevice*)calloc(1, sizeof(MetalDevice)); + PULSE_CHECK_ALLOCATION_RETVAL(metal_device, PULSE_NULL_HANDLE); + +#ifdef PULSE_PLAT_MACOS + uint64_t best = 0; + NSArray>* devices = MTLCopyAllDevices(); + for(id candidate in devices) + { + if(MetalIsDeviceForbidden(candidate, forbiden_devices, forbiden_devices_count)) + continue; + uint64_t current = MetalScoreDevice(candidate); + if(current > best) + { + best = current; + metal_device->device = candidate; + } + } +#endif + if(!metal_device->device) + { + #ifdef PULSE_PLAT_MACOS + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(backend)) + PulseLogError(backend, "(Metal) failed to select device, falling back on default device"); + #endif + id device = MTLCreateSystemDefaultDevice(); + if(!MetalIsDeviceForbidden(device, forbiden_devices, forbiden_devices_count)) + metal_device->device = device; + } + + if(!metal_device->device) + { + if(PULSE_IS_BACKEND_LOW_LEVEL_DEBUG(backend)) + PulseLogError(backend, "(Metal) failed to retrieve default device"); + return PULSE_NULL_HANDLE; + } + + metal_device->queue = [metal_device->device newCommandQueue]; + + pulse_device->driver_data = metal_device; + pulse_device->backend = backend; + PULSE_LOAD_DRIVER_DEVICE(Metal); + + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(backend)) + PulseLogInfoFmt(backend, "(Metal) created device from %s", [metal_device->device.name UTF8String]); + return pulse_device; + } +} + +void MetalDestroyDevice(PulseDevice device) +{ + MetalDevice* metal_device = (MetalDevice*)METAL_RETRIEVE_DRIVER_DATA_AS(device, MetalDevice*); + if(PULSE_IS_BACKEND_HIGH_LEVEL_DEBUG(device->backend)) + PulseLogInfoFmt(device->backend, "(Metal) destroyed device created from %s", [metal_device->device.name UTF8String]); + metal_device->queue = nil; + free(metal_device); + free(device); +} diff --git a/Sources/Backends/Metal/MetalFence.h b/Sources/Backends/Metal/MetalFence.h new file mode 100644 index 0000000..308549b --- /dev/null +++ b/Sources/Backends/Metal/MetalFence.h @@ -0,0 +1,27 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_FENCE_H_ +#define PULSE_METAL_FENCE_H_ + +#include +#include "Metal.h" + +typedef struct MetalFence +{ + int dummy; +} MetalFence; + +PulseFence MetalCreateFence(PulseDevice device); +void MetalDestroyFence(PulseDevice device, PulseFence fence); +bool MetalIsFenceReady(PulseDevice device, PulseFence fence); +bool MetalWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all); + +#endif // PULSE_METAL_FENCE_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalFence.m b/Sources/Backends/Metal/MetalFence.m new file mode 100644 index 0000000..9943fc4 --- /dev/null +++ b/Sources/Backends/Metal/MetalFence.m @@ -0,0 +1,31 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalFence.h" +#include "MetalCommandList.h" + +PulseFence MetalCreateFence(PulseDevice device) +{ + PulseFence fence = (PulseFence)calloc(1, sizeof(PulseFence)); + PULSE_CHECK_ALLOCATION_RETVAL(fence, PULSE_NULL_HANDLE); + return fence; +} + +void MetalDestroyFence(PulseDevice device, PulseFence fence) +{ + free(fence); +} + +bool MetalIsFenceReady(PulseDevice device, PulseFence fence) +{ + return true; +} + +bool MetalWaitForFences(PulseDevice device, const PulseFence* fences, uint32_t fences_count, bool wait_for_all) +{ + return true; +} diff --git a/Sources/Backends/Metal/MetalImage.h b/Sources/Backends/Metal/MetalImage.h new file mode 100644 index 0000000..83e9fb1 --- /dev/null +++ b/Sources/Backends/Metal/MetalImage.h @@ -0,0 +1,27 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include + +#ifdef PULSE_ENABLE_METAL_BACKEND + +#ifndef PULSE_METAL_IMAGE_H_ +#define PULSE_METAL_IMAGE_H_ + +#include "Metal.h" + +typedef struct MetalImage +{ + int dummy; +} MetalImage; + +PulseImage MetalCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos); +bool MetalIsImageFormatValid(PulseDevice device, PulseImageFormat format, PulseImageType type, PulseImageUsageFlags usage); +bool MetalCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, const PulseBufferRegion* dst); +bool MetalBlitImage(PulseCommandList cmd, const PulseImageRegion* src, const PulseImageRegion* dst); +void MetalDestroyImage(PulseDevice device, PulseImage image); + +#endif // PULSE_METAL_IMAGE_H_ + +#endif // PULSE_ENABLE_METAL_BACKEND diff --git a/Sources/Backends/Metal/MetalImage.m b/Sources/Backends/Metal/MetalImage.m new file mode 100644 index 0000000..0c0ae57 --- /dev/null +++ b/Sources/Backends/Metal/MetalImage.m @@ -0,0 +1,28 @@ +// Copyright (C) 2025 kanel +// This file is part of "Pulse" +// For conditions of distribution and use, see copyright notice in LICENSE + +#include +#include "../../PulseInternal.h" +#include "Metal.h" +#include "MetalImage.h" + +PulseImage MetalCreateImage(PulseDevice device, const PulseImageCreateInfo* create_infos) +{ +} + +bool MetalIsImageFormatValid(PulseDevice device, PulseImageFormat format, PulseImageType type, PulseImageUsageFlags usage) +{ +} + +bool MetalCopyImageToBuffer(PulseCommandList cmd, const PulseImageRegion* src, const PulseBufferRegion* dst) +{ +} + +bool MetalBlitImage(PulseCommandList cmd, const PulseImageRegion* src, const PulseImageRegion* dst) +{ +} + +void MetalDestroyImage(PulseDevice device, PulseImage image) +{ +} diff --git a/Sources/Backends/Vulkan/VulkanDevice.c b/Sources/Backends/Vulkan/VulkanDevice.c index 08c1a55..2a62829 100644 --- a/Sources/Backends/Vulkan/VulkanDevice.c +++ b/Sources/Backends/Vulkan/VulkanDevice.c @@ -167,13 +167,17 @@ PulseDevice VulkanCreateDevice(PulseBackend backend, PulseDevice* forbiden_devic instance->vkGetPhysicalDeviceMemoryProperties(device->physical, &device->memory_properties); instance->vkGetPhysicalDeviceFeatures(device->physical, &device->features); + const char* extensions[] = { + "VK_KHR_portability_subset", + }; + VkDeviceCreateInfo create_info = { 0 }; create_info.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO; create_info.queueCreateInfoCount = unique_queues_count; create_info.pQueueCreateInfos = queue_create_infos; create_info.pEnabledFeatures = &device->features; - create_info.enabledExtensionCount = 0; - create_info.ppEnabledExtensionNames = PULSE_NULLPTR; + create_info.enabledExtensionCount = 1; + create_info.ppEnabledExtensionNames = extensions; create_info.enabledLayerCount = 0; create_info.ppEnabledLayerNames = PULSE_NULLPTR; create_info.flags = 0; diff --git a/Sources/Backends/Vulkan/VulkanInstance.c b/Sources/Backends/Vulkan/VulkanInstance.c index fca6bf5..fa49db7 100644 --- a/Sources/Backends/Vulkan/VulkanInstance.c +++ b/Sources/Backends/Vulkan/VulkanInstance.c @@ -115,7 +115,7 @@ static VkInstance VulkanCreateInstance(PulseBackend backend, const char** extens create_info.pApplicationInfo = &app_info; create_info.enabledExtensionCount = extensions_count; create_info.ppEnabledExtensionNames = extensions_enabled; - #ifdef PULSE_PLAT_MACOS + #ifdef PULSE_PLAT_APPLE create_info.flags = VK_INSTANCE_CREATE_ENUMERATE_PORTABILITY_BIT_KHR; #else create_info.flags = 0; @@ -151,7 +151,7 @@ static VkInstance VulkanCreateInstance(PulseBackend backend, const char** extens bool VulkanInitInstance(PulseBackend backend, VulkanInstance* instance, PulseDebugLevel debug_level) { - #ifdef PULSE_PLAT_MACOS + #ifdef PULSE_PLAT_APPLE const char* extensions[] = { VK_KHR_PORTABILITY_ENUMERATION_EXTENSION_NAME, }; diff --git a/Sources/Backends/Vulkan/VulkanLoader.c b/Sources/Backends/Vulkan/VulkanLoader.c index 2b5f675..3865e96 100644 --- a/Sources/Backends/Vulkan/VulkanLoader.c +++ b/Sources/Backends/Vulkan/VulkanLoader.c @@ -40,7 +40,7 @@ bool VulkanInitLoader() const char* libnames[] = { "vulkan-1.dll" }; - #elif defined(PULSE_PLAT_MACOS) + #elif defined(PULSE_PLAT_APPLE) const char* libnames[] = { "libvulkan.dylib", "libvulkan.1.dylib", diff --git a/Sources/PulseDefs.h b/Sources/PulseDefs.h index b2dcfca..db84664 100644 --- a/Sources/PulseDefs.h +++ b/Sources/PulseDefs.h @@ -101,6 +101,8 @@ #endif #endif +#define PULSE_ALIGN_UP(val, alignment) ((val + alignment - 1) & ~(alignment - 1)) + #define PULSE_LOAD_DRIVER_DEVICE_FUNCTION(fn, _namespace) pulse_device->PFN_##fn = _namespace##fn; #define PULSE_LOAD_DRIVER_DEVICE(_namespace) \ PULSE_LOAD_DRIVER_DEVICE_FUNCTION(DestroyDevice, _namespace) \ diff --git a/Sources/PulseInternal.c b/Sources/PulseInternal.c index 5a3b12e..c031ae6 100644 --- a/Sources/PulseInternal.c +++ b/Sources/PulseInternal.c @@ -42,7 +42,7 @@ PULSE_IMPORT_API FARPROC __stdcall GetProcAddress(HMODULE, LPCSTR); PULSE_IMPORT_API int __stdcall FreeLibrary(HMODULE); #else - #ifdef PULSE_PLAT_MACOS + #ifdef PULSE_PLAT_APPLE #include #endif #include diff --git a/xmake.lua b/xmake.lua index 5f75a2c..2cfc834 100644 --- a/xmake.lua +++ b/xmake.lua @@ -60,7 +60,7 @@ local backends = { end, before_build = function(target, os) local gles_dir = target:pkg("opengl-headers"):installdir() - os.runv("python", {"Scripts/GenerateOpenGLDefs.py", "Sources/Backends/OpenGL/OpenGLFunctions.h", gles_dir .. "/include/GLES3/gl32.h", "Sources/Backends/OpenGL/OpenGLWraps.h"}) + os.runv("python3", {"Scripts/GenerateOpenGLDefs.py", "Sources/Backends/OpenGL/OpenGLFunctions.h", gles_dir .. "/include/GLES3/gl32.h", "Sources/Backends/OpenGL/OpenGLWraps.h"}) end }, } @@ -157,7 +157,11 @@ target("pulse_gpu") add_headerfiles("Sources/Backends/" .. name .. "/**.h", { prefixdir = "private", install = false }) add_headerfiles("Sources/Backends/" .. name .. "/**.inl", { prefixdir = "private", install = false }) - add_files("Sources/Backends/" .. name .. "/**.c") + -- Checks if there are C files in the backend directory and add them if so + local cfiles = os.files("Sources/Backends/" .. name .. "/**.c") + if #cfiles > 0 then + add_files("Sources/Backends/" .. name .. "/**.c") + end if module.custom then module.custom()