adding device selection, buffer creation and pipeline creation to Metal backend

This commit is contained in:
2025-09-07 01:04:15 +02:00
parent 231e8f2687
commit e9f01cf268
27 changed files with 848 additions and 16 deletions

86
.github/workflows/metal-test-macos.yml vendored git.filemode.normal_file
View File

@@ -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

63
Examples/Metal/main.c git.filemode.normal_file
View File

@@ -0,0 +1,63 @@
#include <Pulse.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
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 <metal_stdlib>\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<int>(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;
}

4
Examples/Metal/xmake.lua git.filemode.normal_file
View File

@@ -0,0 +1,4 @@
target("MetalExample")
add_deps("pulse_gpu")
add_files("*.c")
target_end()

View File

@@ -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

View File

@@ -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 <TargetConditionals.h>
#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

View File

@@ -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_

View File

@@ -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

29
Sources/Backends/Metal/MetalBuffer.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#ifdef PULSE_ENABLE_METAL_BACKEND
#ifndef PULSE_METAL_BUFFER_H_
#define PULSE_METAL_BUFFER_H_
#include <Metal/Metal.h>
#include "Metal.h"
typedef struct MetalBuffer
{
id<MTLBuffer> 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

73
Sources/Backends/Metal/MetalBuffer.m git.filemode.normal_file
View File

@@ -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 <string.h>
#include <Pulse.h>
#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);
}

29
Sources/Backends/Metal/MetalCommandList.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#ifdef PULSE_ENABLE_METAL_BACKEND
#ifndef PULSE_METAL_COMMAND_LIST_H_
#define PULSE_METAL_COMMAND_LIST_H_
#include <stdatomic.h>
#include <tinycthread.h>
#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

29
Sources/Backends/Metal/MetalCommandList.m git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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)
{
}

27
Sources/Backends/Metal/MetalComputePass.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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

49
Sources/Backends/Metal/MetalComputePass.m git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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)
{
}

25
Sources/Backends/Metal/MetalComputePipeline.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#ifdef PULSE_ENABLE_METAL_BACKEND
#ifndef PULSE_METAL_COMPUTE_PIPELINE_H_
#define PULSE_METAL_COMPUTE_PIPELINE_H_
#include <Metal/Metal.h>
#include "Metal.h"
typedef struct MetalComputePipeline
{
id<MTLComputePipelineState> 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

109
Sources/Backends/Metal/MetalComputePipeline.m git.filemode.normal_file
View File

@@ -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 <string.h>
#include <Pulse.h>
#include "../../PulseInternal.h"
#include "Metal.h"
#include "MetalDevice.h"
#include "MetalComputePipeline.h"
typedef struct MetalLibraryFunction
{
id<MTLLibrary> library;
id<MTLFunction> 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<MTLLibrary> 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<MTLFunction> 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);
}

26
Sources/Backends/Metal/MetalDevice.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#ifdef PULSE_ENABLE_METAL_BACKEND
#ifndef PULSE_METAL_DEVICE_H_
#define PULSE_METAL_DEVICE_H_
#include <Metal/Metal.h>
#include "Metal.h"
typedef struct MetalDevice
{
id<MTLDevice> device;
id<MTLCommandQueue> 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

115
Sources/Backends/Metal/MetalDevice.m git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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<MTLDevice> 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<MTLDevice> 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<id<MTLDevice>>* devices = MTLCopyAllDevices();
for(id<MTLDevice> 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<MTLDevice> 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);
}

27
Sources/Backends/Metal/MetalFence.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#ifdef PULSE_ENABLE_METAL_BACKEND
#ifndef PULSE_METAL_FENCE_H_
#define PULSE_METAL_FENCE_H_
#include <Pulse.h>
#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

31
Sources/Backends/Metal/MetalFence.m git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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;
}

27
Sources/Backends/Metal/MetalImage.h git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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

28
Sources/Backends/Metal/MetalImage.m git.filemode.normal_file
View File

@@ -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 <Pulse.h>
#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)
{
}

View File

@@ -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;

View File

@@ -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,
};

View File

@@ -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",

View File

@@ -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) \

View File

@@ -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 <stdlib.h>
#endif
#include <dlfcn.h>

View File

@@ -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()