diff --git a/src/kosmickrisp/bridge/mtl_device.h b/src/kosmickrisp/bridge/mtl_device.h index c180c9446e3..1732f924750 100644 --- a/src/kosmickrisp/bridge/mtl_device.h +++ b/src/kosmickrisp/bridge/mtl_device.h @@ -28,6 +28,8 @@ uint64_t mtl_device_get_peer_group_id(mtl_device *dev); uint32_t mtl_device_get_peer_index(mtl_device *dev); uint64_t mtl_device_get_registry_id(mtl_device *dev); struct mtl_size mtl_device_max_threads_per_threadgroup(mtl_device *dev); +uint32_t mtl_device_max_threadgroup_memory_length(mtl_device *dev); +uint64_t mtl_device_max_buffer_length(mtl_device *dev); /* Timestamp query */ uint64_t mtl_device_get_gpu_timestamp(mtl_device *dev); diff --git a/src/kosmickrisp/bridge/mtl_device.m b/src/kosmickrisp/bridge/mtl_device.m index 33aec0987da..6159d95fb31 100644 --- a/src/kosmickrisp/bridge/mtl_device.m +++ b/src/kosmickrisp/bridge/mtl_device.m @@ -129,6 +129,24 @@ mtl_device_max_threads_per_threadgroup(mtl_device *dev) } } +uint32_t +mtl_device_max_threadgroup_memory_length(mtl_device *dev) +{ + @autoreleasepool { + id device = (id)dev; + return device.maxThreadgroupMemoryLength; + } +} + +uint64_t +mtl_device_max_buffer_length(mtl_device *dev) +{ + @autoreleasepool { + id device = (id)dev; + return device.maxBufferLength; + } +} + /* Timestamp query */ uint64_t mtl_device_get_gpu_timestamp(mtl_device *dev) diff --git a/src/kosmickrisp/bridge/stubs/mtl_device.c b/src/kosmickrisp/bridge/stubs/mtl_device.c index 260da172de2..26d19cc1d50 100644 --- a/src/kosmickrisp/bridge/stubs/mtl_device.c +++ b/src/kosmickrisp/bridge/stubs/mtl_device.c @@ -59,6 +59,18 @@ mtl_device_max_threads_per_threadgroup(mtl_device *dev) return (struct mtl_size){}; } +uint32_t +mtl_device_max_threadgroup_memory_length(mtl_device *dev) +{ + return 0u; +} + +uint64_t +mtl_device_max_buffer_length(mtl_device *dev) +{ + return 0u; +} + /* Timestamp query */ uint64_t mtl_device_get_gpu_timestamp(mtl_device *dev) diff --git a/src/kosmickrisp/vulkan/kk_buffer.c b/src/kosmickrisp/vulkan/kk_buffer.c index 75c10ab86f9..56c6ea3337f 100644 --- a/src/kosmickrisp/vulkan/kk_buffer.c +++ b/src/kosmickrisp/vulkan/kk_buffer.c @@ -35,9 +35,10 @@ kk_CreateBuffer(VkDevice device, const VkBufferCreateInfo *pCreateInfo, const VkAllocationCallbacks *pAllocator, VkBuffer *pBuffer) { VK_FROM_HANDLE(kk_device, dev, device); + struct kk_physical_device *pdev = kk_device_physical(dev); struct kk_buffer *buffer; - if (pCreateInfo->size > KK_MAX_BUFFER_SIZE) + if (pCreateInfo->size > pdev->info.max_buffer_size) return vk_error(dev, VK_ERROR_OUT_OF_DEVICE_MEMORY); buffer = diff --git a/src/kosmickrisp/vulkan/kk_physical_device.c b/src/kosmickrisp/vulkan/kk_physical_device.c index 6dbfc538420..bf146e0bbd3 100644 --- a/src/kosmickrisp/vulkan/kk_physical_device.c +++ b/src/kosmickrisp/vulkan/kk_physical_device.c @@ -400,7 +400,7 @@ kk_get_device_properties(const struct kk_physical_device *pdev, .maxFragmentOutputAttachments = KK_MAX_RTS, .maxFragmentDualSrcAttachments = 1, .maxFragmentCombinedOutputResources = KK_MAX_DESCRIPTORS, - .maxComputeSharedMemorySize = KK_MAX_SHARED_SIZE, + .maxComputeSharedMemorySize = pdev->info.max_compute_shared_memory_size, .maxComputeWorkGroupCount = {0x7fffffff, 65535, 65535}, .maxComputeWorkGroupInvocations = pdev->info.max_workgroup_invocations, .maxComputeWorkGroupSize = {pdev->info.max_workgroup_count[0], @@ -483,7 +483,7 @@ kk_get_device_properties(const struct kk_physical_device *pdev, .maxMultiviewViewCount = KK_MAX_MULTIVIEW_VIEW_COUNT, .maxMultiviewInstanceIndex = UINT32_MAX, .maxPerSetDescriptors = UINT32_MAX, - .maxMemoryAllocationSize = (1u << 31), + .maxMemoryAllocationSize = pdev->info.max_buffer_size, /* Vulkan 1.2 properties */ .supportedDepthResolveModes = @@ -562,7 +562,7 @@ kk_get_device_properties(const struct kk_physical_device *pdev, .storageTexelBufferOffsetSingleTexelAlignment = false, .uniformTexelBufferOffsetAlignmentBytes = KK_MIN_TEXEL_BUFFER_ALIGNMENT, .uniformTexelBufferOffsetSingleTexelAlignment = false, - .maxBufferSize = KK_MAX_BUFFER_SIZE, + .maxBufferSize = pdev->info.max_buffer_size, /* VK_KHR_push_descriptor */ .maxPushDescriptors = KK_MAX_PUSH_DESCRIPTORS, @@ -798,6 +798,11 @@ get_metal_limits(struct kk_physical_device *pdev) pdev->info.max_workgroup_count[2] = workgroup_size.z; pdev->info.max_workgroup_invocations = MAX3(workgroup_size.x, workgroup_size.y, workgroup_size.z); + + pdev->info.max_compute_shared_memory_size = + mtl_device_max_threadgroup_memory_length(pdev->mtl_dev_handle); + pdev->info.max_buffer_size = + mtl_device_max_buffer_length(pdev->mtl_dev_handle); } VkResult diff --git a/src/kosmickrisp/vulkan/kk_physical_device.h b/src/kosmickrisp/vulkan/kk_physical_device.h index 9daf6dfbdcc..91c053a8310 100644 --- a/src/kosmickrisp/vulkan/kk_physical_device.h +++ b/src/kosmickrisp/vulkan/kk_physical_device.h @@ -38,6 +38,8 @@ struct kk_memory_heap { struct kk_device_info { uint32_t max_workgroup_count[3]; uint32_t max_workgroup_invocations; + uint32_t max_compute_shared_memory_size; + uint64_t max_buffer_size; }; struct kk_physical_device { diff --git a/src/kosmickrisp/vulkan/kk_private.h b/src/kosmickrisp/vulkan/kk_private.h index edfa78351c8..efdf7d6acc1 100644 --- a/src/kosmickrisp/vulkan/kk_private.h +++ b/src/kosmickrisp/vulkan/kk_private.h @@ -37,8 +37,6 @@ #define KK_MAX_OCCLUSION_QUERIES (32768) #define KK_SPARSE_ADDR_SPACE_SIZE (1ull << 39) -#define KK_MAX_BUFFER_SIZE (1ull << 31) -#define KK_MAX_SHARED_SIZE (32 * 1024) /* Max size of a bound cbuf */ #define KK_MAX_CBUF_SIZE (1u << 16)