kk: Use device limits for buffers and compute shared memory.

Metal provides these limits as properties of MTLDevice, which can be
used instead of hardcoding them.

Reviewed-by: Aitor Camacho <aitor@lunarg.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41077>

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

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<MTLDevice> device = (id<MTLDevice>)dev;
+      return device.maxThreadgroupMemoryLength;
+   }
+}
+
+uint64_t
+mtl_device_max_buffer_length(mtl_device *dev)
+{
+   @autoreleasepool {
+      id<MTLDevice> device = (id<MTLDevice>)dev;
+      return device.maxBufferLength;
+   }
+}
+
 /* Timestamp query */
 uint64_t
 mtl_device_get_gpu_timestamp(mtl_device *dev)

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)

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 =

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

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 {

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)