1 /*
2 * Copyright © 2022 Konstantin Seurer
3 *
4 * SPDX-License-Identifier: MIT
5 */
6
7 #include "radv_radix_sort.h"
8 #include "targets/u64/config.h"
9 #include "radv_cmd_buffer.h"
10 #include "target.h"
11
12 static const uint32_t init_spv[] = {
13 #include "radix_sort/shaders/init.comp.spv.h"
14 };
15
16 static const uint32_t fill_spv[] = {
17 #include "radix_sort/shaders/fill.comp.spv.h"
18 };
19
20 static const uint32_t histogram_spv[] = {
21 #include "radix_sort/shaders/histogram.comp.spv.h"
22 };
23
24 static const uint32_t prefix_spv[] = {
25 #include "radix_sort/shaders/prefix.comp.spv.h"
26 };
27
28 static const uint32_t scatter_0_even_spv[] = {
29 #include "radix_sort/shaders/scatter_0_even.comp.spv.h"
30 };
31
32 static const uint32_t scatter_0_odd_spv[] = {
33 #include "radix_sort/shaders/scatter_0_odd.comp.spv.h"
34 };
35
36 static const uint32_t scatter_1_even_spv[] = {
37 #include "radix_sort/shaders/scatter_1_even.comp.spv.h"
38 };
39
40 static const uint32_t scatter_1_odd_spv[] = {
41 #include "radix_sort/shaders/scatter_1_odd.comp.spv.h"
42 };
43
44 static const struct radix_sort_vk_target_config target_config = {
45 .keyval_dwords = RS_KEYVAL_DWORDS,
46
47 .histogram =
48 {
49 .workgroup_size_log2 = RS_HISTOGRAM_WORKGROUP_SIZE_LOG2,
50 .subgroup_size_log2 = RS_HISTOGRAM_SUBGROUP_SIZE_LOG2,
51 .block_rows = RS_HISTOGRAM_BLOCK_ROWS,
52 },
53
54 .prefix =
55 {
56 .workgroup_size_log2 = RS_PREFIX_WORKGROUP_SIZE_LOG2,
57 .subgroup_size_log2 = RS_PREFIX_SUBGROUP_SIZE_LOG2,
58 },
59
60 .scatter =
61 {
62 .workgroup_size_log2 = RS_SCATTER_WORKGROUP_SIZE_LOG2,
63 .subgroup_size_log2 = RS_SCATTER_SUBGROUP_SIZE_LOG2,
64 .block_rows = RS_SCATTER_BLOCK_ROWS,
65 },
66 };
67
68 radix_sort_vk_t *
radv_create_radix_sort_u64(VkDevice device,VkAllocationCallbacks const * ac,VkPipelineCache pc)69 radv_create_radix_sort_u64(VkDevice device, VkAllocationCallbacks const *ac, VkPipelineCache pc)
70 {
71 const uint32_t *spv[8] = {
72 init_spv, fill_spv, histogram_spv, prefix_spv,
73 scatter_0_even_spv, scatter_0_odd_spv, scatter_1_even_spv, scatter_1_odd_spv,
74 };
75 const uint32_t spv_sizes[8] = {
76 sizeof(init_spv), sizeof(fill_spv), sizeof(histogram_spv), sizeof(prefix_spv),
77 sizeof(scatter_0_even_spv), sizeof(scatter_0_odd_spv), sizeof(scatter_1_even_spv), sizeof(scatter_1_odd_spv),
78 };
79 return radix_sort_vk_create(device, ac, pc, spv, spv_sizes, target_config);
80 }
81
82 VKAPI_ATTR VkResult VKAPI_CALL
vkCreateShaderModule(VkDevice _device,const VkShaderModuleCreateInfo * pCreateInfo,const VkAllocationCallbacks * pAllocator,VkShaderModule * pShaderModule)83 vkCreateShaderModule(VkDevice _device, const VkShaderModuleCreateInfo *pCreateInfo,
84 const VkAllocationCallbacks *pAllocator, VkShaderModule *pShaderModule)
85 {
86 VK_FROM_HANDLE(radv_device, device, _device);
87 return device->vk.dispatch_table.CreateShaderModule(_device, pCreateInfo, pAllocator, pShaderModule);
88 }
89
90 VKAPI_ATTR void VKAPI_CALL
vkDestroyShaderModule(VkDevice _device,VkShaderModule shaderModule,const VkAllocationCallbacks * pAllocator)91 vkDestroyShaderModule(VkDevice _device, VkShaderModule shaderModule, const VkAllocationCallbacks *pAllocator)
92 {
93 VK_FROM_HANDLE(radv_device, device, _device);
94 device->vk.dispatch_table.DestroyShaderModule(_device, shaderModule, pAllocator);
95 }
96
97 VKAPI_ATTR VkResult VKAPI_CALL
vkCreatePipelineLayout(VkDevice _device,const VkPipelineLayoutCreateInfo * pCreateInfo,const VkAllocationCallbacks * pAllocator,VkPipelineLayout * pPipelineLayout)98 vkCreatePipelineLayout(VkDevice _device, const VkPipelineLayoutCreateInfo *pCreateInfo,
99 const VkAllocationCallbacks *pAllocator, VkPipelineLayout *pPipelineLayout)
100 {
101 VK_FROM_HANDLE(radv_device, device, _device);
102 return device->vk.dispatch_table.CreatePipelineLayout(_device, pCreateInfo, pAllocator, pPipelineLayout);
103 }
104
105 VKAPI_ATTR void VKAPI_CALL
vkDestroyPipelineLayout(VkDevice _device,VkPipelineLayout pipelineLayout,const VkAllocationCallbacks * pAllocator)106 vkDestroyPipelineLayout(VkDevice _device, VkPipelineLayout pipelineLayout, const VkAllocationCallbacks *pAllocator)
107 {
108 VK_FROM_HANDLE(radv_device, device, _device);
109 device->vk.dispatch_table.DestroyPipelineLayout(_device, pipelineLayout, pAllocator);
110 }
111
112 VKAPI_ATTR VkResult VKAPI_CALL
vkCreateComputePipelines(VkDevice _device,VkPipelineCache pipelineCache,uint32_t createInfoCount,const VkComputePipelineCreateInfo * pCreateInfos,const VkAllocationCallbacks * pAllocator,VkPipeline * pPipelines)113 vkCreateComputePipelines(VkDevice _device, VkPipelineCache pipelineCache, uint32_t createInfoCount,
114 const VkComputePipelineCreateInfo *pCreateInfos, const VkAllocationCallbacks *pAllocator,
115 VkPipeline *pPipelines)
116 {
117 VK_FROM_HANDLE(radv_device, device, _device);
118 return device->vk.dispatch_table.CreateComputePipelines(_device, pipelineCache, createInfoCount, pCreateInfos,
119 pAllocator, pPipelines);
120 }
121
122 VKAPI_ATTR void VKAPI_CALL
vkDestroyPipeline(VkDevice _device,VkPipeline pipeline,const VkAllocationCallbacks * pAllocator)123 vkDestroyPipeline(VkDevice _device, VkPipeline pipeline, const VkAllocationCallbacks *pAllocator)
124 {
125 VK_FROM_HANDLE(radv_device, device, _device);
126 device->vk.dispatch_table.DestroyPipeline(_device, pipeline, pAllocator);
127 }
128
129 VKAPI_ATTR void VKAPI_CALL
vkCmdPipelineBarrier(VkCommandBuffer commandBuffer,VkPipelineStageFlags srcStageMask,VkPipelineStageFlags dstStageMask,VkDependencyFlags dependencyFlags,uint32_t memoryBarrierCount,const VkMemoryBarrier * pMemoryBarriers,uint32_t bufferMemoryBarrierCount,const VkBufferMemoryBarrier * pBufferMemoryBarriers,uint32_t imageMemoryBarrierCount,const VkImageMemoryBarrier * pImageMemoryBarriers)130 vkCmdPipelineBarrier(VkCommandBuffer commandBuffer, VkPipelineStageFlags srcStageMask,
131 VkPipelineStageFlags dstStageMask, VkDependencyFlags dependencyFlags, uint32_t memoryBarrierCount,
132 const VkMemoryBarrier *pMemoryBarriers, uint32_t bufferMemoryBarrierCount,
133 const VkBufferMemoryBarrier *pBufferMemoryBarriers, uint32_t imageMemoryBarrierCount,
134 const VkImageMemoryBarrier *pImageMemoryBarriers)
135 {
136 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
137 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
138
139 device->vk.dispatch_table.CmdPipelineBarrier(commandBuffer, srcStageMask, dstStageMask, dependencyFlags,
140 memoryBarrierCount, pMemoryBarriers, bufferMemoryBarrierCount,
141 pBufferMemoryBarriers, imageMemoryBarrierCount, pImageMemoryBarriers);
142 }
143
144 VKAPI_ATTR void VKAPI_CALL
vkCmdPushConstants(VkCommandBuffer commandBuffer,VkPipelineLayout layout,VkShaderStageFlags stageFlags,uint32_t offset,uint32_t size,const void * pValues)145 vkCmdPushConstants(VkCommandBuffer commandBuffer, VkPipelineLayout layout, VkShaderStageFlags stageFlags,
146 uint32_t offset, uint32_t size, const void *pValues)
147 {
148 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
149 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
150
151 device->vk.dispatch_table.CmdPushConstants(commandBuffer, layout, stageFlags, offset, size, pValues);
152 }
153
154 VKAPI_ATTR void VKAPI_CALL
vkCmdBindPipeline(VkCommandBuffer commandBuffer,VkPipelineBindPoint pipelineBindPoint,VkPipeline pipeline)155 vkCmdBindPipeline(VkCommandBuffer commandBuffer, VkPipelineBindPoint pipelineBindPoint, VkPipeline pipeline)
156 {
157 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
158 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
159
160 device->vk.dispatch_table.CmdBindPipeline(commandBuffer, pipelineBindPoint, pipeline);
161 }
162
163 VKAPI_ATTR void VKAPI_CALL
vkCmdDispatch(VkCommandBuffer commandBuffer,uint32_t groupCountX,uint32_t groupCountY,uint32_t groupCountZ)164 vkCmdDispatch(VkCommandBuffer commandBuffer, uint32_t groupCountX, uint32_t groupCountY, uint32_t groupCountZ)
165 {
166 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
167 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
168
169 device->vk.dispatch_table.CmdDispatch(commandBuffer, groupCountX, groupCountY, groupCountZ);
170 }
171
172 VKAPI_ATTR VkDeviceAddress VKAPI_CALL
vkGetBufferDeviceAddress(VkDevice _device,const VkBufferDeviceAddressInfo * pInfo)173 vkGetBufferDeviceAddress(VkDevice _device, const VkBufferDeviceAddressInfo *pInfo)
174 {
175 VK_FROM_HANDLE(radv_device, device, _device);
176 return device->vk.dispatch_table.GetBufferDeviceAddress(_device, pInfo);
177 }
178
179 VKAPI_ATTR void VKAPI_CALL
vkCmdFillBuffer(VkCommandBuffer commandBuffer,VkBuffer dstBuffer,VkDeviceSize dstOffset,VkDeviceSize size,uint32_t data)180 vkCmdFillBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer, VkDeviceSize dstOffset, VkDeviceSize size,
181 uint32_t data)
182 {
183 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
184 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
185
186 device->vk.dispatch_table.CmdFillBuffer(commandBuffer, dstBuffer, dstOffset, size, data);
187 }
188
189 VKAPI_ATTR void VKAPI_CALL
vkCmdDispatchIndirect(VkCommandBuffer commandBuffer,VkBuffer buffer,VkDeviceSize offset)190 vkCmdDispatchIndirect(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize offset)
191 {
192 VK_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
193 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
194
195 device->vk.dispatch_table.CmdDispatchIndirect(commandBuffer, buffer, offset);
196 }
197