1 /*
2 * Copyright © 2022 Konstantin Seurer
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 */
23
24 #include "radv_radix_sort.h"
25 #include "targets/u64/config.h"
26 #include "radv_private.h"
27 #include "target.h"
28
29 static const uint32_t init_spv[] = {
30 #include "radix_sort/shaders/init.comp.spv.h"
31 };
32
33 static const uint32_t fill_spv[] = {
34 #include "radix_sort/shaders/fill.comp.spv.h"
35 };
36
37 static const uint32_t histogram_spv[] = {
38 #include "radix_sort/shaders/histogram.comp.spv.h"
39 };
40
41 static const uint32_t prefix_spv[] = {
42 #include "radix_sort/shaders/prefix.comp.spv.h"
43 };
44
45 static const uint32_t scatter_0_even_spv[] = {
46 #include "radix_sort/shaders/scatter_0_even.comp.spv.h"
47 };
48
49 static const uint32_t scatter_0_odd_spv[] = {
50 #include "radix_sort/shaders/scatter_0_odd.comp.spv.h"
51 };
52
53 static const uint32_t scatter_1_even_spv[] = {
54 #include "radix_sort/shaders/scatter_1_even.comp.spv.h"
55 };
56
57 static const uint32_t scatter_1_odd_spv[] = {
58 #include "radix_sort/shaders/scatter_1_odd.comp.spv.h"
59 };
60
61 static const struct radix_sort_vk_target_config target_config = {
62 .keyval_dwords = RS_KEYVAL_DWORDS,
63
64 .histogram =
65 {
66 .workgroup_size_log2 = RS_HISTOGRAM_WORKGROUP_SIZE_LOG2,
67 .subgroup_size_log2 = RS_HISTOGRAM_SUBGROUP_SIZE_LOG2,
68 .block_rows = RS_HISTOGRAM_BLOCK_ROWS,
69 },
70
71 .prefix =
72 {
73 .workgroup_size_log2 = RS_PREFIX_WORKGROUP_SIZE_LOG2,
74 .subgroup_size_log2 = RS_PREFIX_SUBGROUP_SIZE_LOG2,
75 },
76
77 .scatter =
78 {
79 .workgroup_size_log2 = RS_SCATTER_WORKGROUP_SIZE_LOG2,
80 .subgroup_size_log2 = RS_SCATTER_SUBGROUP_SIZE_LOG2,
81 .block_rows = RS_SCATTER_BLOCK_ROWS,
82 },
83 };
84
85 radix_sort_vk_t *
radv_create_radix_sort_u64(VkDevice device,VkAllocationCallbacks const * ac,VkPipelineCache pc)86 radv_create_radix_sort_u64(VkDevice device, VkAllocationCallbacks const *ac, VkPipelineCache pc)
87 {
88 const uint32_t *spv[8] = {
89 init_spv, fill_spv, histogram_spv, prefix_spv,
90 scatter_0_even_spv, scatter_0_odd_spv, scatter_1_even_spv, scatter_1_odd_spv,
91 };
92 const uint32_t spv_sizes[8] = {
93 sizeof(init_spv), sizeof(fill_spv), sizeof(histogram_spv),
94 sizeof(prefix_spv), sizeof(scatter_0_even_spv), sizeof(scatter_0_odd_spv),
95 sizeof(scatter_1_even_spv), sizeof(scatter_1_odd_spv),
96 };
97 return radix_sort_vk_create(device, ac, pc, spv, spv_sizes, target_config);
98 }
99
100 VKAPI_ATTR VkResult VKAPI_CALL
vkCreateShaderModule(VkDevice _device,const VkShaderModuleCreateInfo * pCreateInfo,const VkAllocationCallbacks * pAllocator,VkShaderModule * pShaderModule)101 vkCreateShaderModule(VkDevice _device, const VkShaderModuleCreateInfo *pCreateInfo,
102 const VkAllocationCallbacks *pAllocator, VkShaderModule *pShaderModule)
103 {
104 RADV_FROM_HANDLE(radv_device, device, _device);
105 return device->vk.dispatch_table.CreateShaderModule(_device, pCreateInfo, pAllocator,
106 pShaderModule);
107 }
108
109 VKAPI_ATTR void VKAPI_CALL
vkDestroyShaderModule(VkDevice _device,VkShaderModule shaderModule,const VkAllocationCallbacks * pAllocator)110 vkDestroyShaderModule(VkDevice _device, VkShaderModule shaderModule,
111 const VkAllocationCallbacks *pAllocator)
112 {
113 RADV_FROM_HANDLE(radv_device, device, _device);
114 device->vk.dispatch_table.DestroyShaderModule(_device, shaderModule, pAllocator);
115 }
116
117 VKAPI_ATTR VkResult VKAPI_CALL
vkCreatePipelineLayout(VkDevice _device,const VkPipelineLayoutCreateInfo * pCreateInfo,const VkAllocationCallbacks * pAllocator,VkPipelineLayout * pPipelineLayout)118 vkCreatePipelineLayout(VkDevice _device, const VkPipelineLayoutCreateInfo *pCreateInfo,
119 const VkAllocationCallbacks *pAllocator, VkPipelineLayout *pPipelineLayout)
120 {
121 RADV_FROM_HANDLE(radv_device, device, _device);
122 return device->vk.dispatch_table.CreatePipelineLayout(_device, pCreateInfo, pAllocator,
123 pPipelineLayout);
124 }
125
126 VKAPI_ATTR void VKAPI_CALL
vkDestroyPipelineLayout(VkDevice _device,VkPipelineLayout pipelineLayout,const VkAllocationCallbacks * pAllocator)127 vkDestroyPipelineLayout(VkDevice _device, VkPipelineLayout pipelineLayout,
128 const VkAllocationCallbacks *pAllocator)
129 {
130 RADV_FROM_HANDLE(radv_device, device, _device);
131 device->vk.dispatch_table.DestroyPipelineLayout(_device, pipelineLayout, pAllocator);
132 }
133
134 VKAPI_ATTR VkResult VKAPI_CALL
vkCreateComputePipelines(VkDevice _device,VkPipelineCache pipelineCache,uint32_t createInfoCount,const VkComputePipelineCreateInfo * pCreateInfos,const VkAllocationCallbacks * pAllocator,VkPipeline * pPipelines)135 vkCreateComputePipelines(VkDevice _device, VkPipelineCache pipelineCache, uint32_t createInfoCount,
136 const VkComputePipelineCreateInfo *pCreateInfos,
137 const VkAllocationCallbacks *pAllocator, VkPipeline *pPipelines)
138 {
139 RADV_FROM_HANDLE(radv_device, device, _device);
140 return device->vk.dispatch_table.CreateComputePipelines(_device, pipelineCache, createInfoCount,
141 pCreateInfos, pAllocator, pPipelines);
142 }
143
144 VKAPI_ATTR void VKAPI_CALL
vkDestroyPipeline(VkDevice _device,VkPipeline pipeline,const VkAllocationCallbacks * pAllocator)145 vkDestroyPipeline(VkDevice _device, VkPipeline pipeline, const VkAllocationCallbacks *pAllocator)
146 {
147 RADV_FROM_HANDLE(radv_device, device, _device);
148 device->vk.dispatch_table.DestroyPipeline(_device, pipeline, pAllocator);
149 }
150
151 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)152 vkCmdPipelineBarrier(VkCommandBuffer commandBuffer, VkPipelineStageFlags srcStageMask,
153 VkPipelineStageFlags dstStageMask, VkDependencyFlags dependencyFlags,
154 uint32_t memoryBarrierCount, const VkMemoryBarrier *pMemoryBarriers,
155 uint32_t bufferMemoryBarrierCount,
156 const VkBufferMemoryBarrier *pBufferMemoryBarriers,
157 uint32_t imageMemoryBarrierCount,
158 const VkImageMemoryBarrier *pImageMemoryBarriers)
159 {
160 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
161 cmd_buffer->device->vk.dispatch_table.CmdPipelineBarrier(
162 commandBuffer, srcStageMask, dstStageMask, dependencyFlags, memoryBarrierCount,
163 pMemoryBarriers, bufferMemoryBarrierCount, pBufferMemoryBarriers, imageMemoryBarrierCount,
164 pImageMemoryBarriers);
165 }
166
167 VKAPI_ATTR void VKAPI_CALL
vkCmdPushConstants(VkCommandBuffer commandBuffer,VkPipelineLayout layout,VkShaderStageFlags stageFlags,uint32_t offset,uint32_t size,const void * pValues)168 vkCmdPushConstants(VkCommandBuffer commandBuffer, VkPipelineLayout layout,
169 VkShaderStageFlags stageFlags, uint32_t offset, uint32_t size,
170 const void *pValues)
171 {
172 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
173 cmd_buffer->device->vk.dispatch_table.CmdPushConstants(commandBuffer, layout, stageFlags, offset,
174 size, pValues);
175 }
176
177 VKAPI_ATTR void VKAPI_CALL
vkCmdBindPipeline(VkCommandBuffer commandBuffer,VkPipelineBindPoint pipelineBindPoint,VkPipeline pipeline)178 vkCmdBindPipeline(VkCommandBuffer commandBuffer, VkPipelineBindPoint pipelineBindPoint,
179 VkPipeline pipeline)
180 {
181 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
182 cmd_buffer->device->vk.dispatch_table.CmdBindPipeline(commandBuffer, pipelineBindPoint,
183 pipeline);
184 }
185
186 VKAPI_ATTR void VKAPI_CALL
vkCmdDispatch(VkCommandBuffer commandBuffer,uint32_t groupCountX,uint32_t groupCountY,uint32_t groupCountZ)187 vkCmdDispatch(VkCommandBuffer commandBuffer, uint32_t groupCountX, uint32_t groupCountY,
188 uint32_t groupCountZ)
189 {
190 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
191 cmd_buffer->device->vk.dispatch_table.CmdDispatch(commandBuffer, groupCountX, groupCountY,
192 groupCountZ);
193 }
194
195 VKAPI_ATTR VkDeviceAddress VKAPI_CALL
vkGetBufferDeviceAddress(VkDevice _device,const VkBufferDeviceAddressInfo * pInfo)196 vkGetBufferDeviceAddress(VkDevice _device, const VkBufferDeviceAddressInfo *pInfo)
197 {
198 RADV_FROM_HANDLE(radv_device, device, _device);
199 return device->vk.dispatch_table.GetBufferDeviceAddress(_device, pInfo);
200 }
201
202 VKAPI_ATTR void VKAPI_CALL
vkCmdFillBuffer(VkCommandBuffer commandBuffer,VkBuffer dstBuffer,VkDeviceSize dstOffset,VkDeviceSize size,uint32_t data)203 vkCmdFillBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer, VkDeviceSize dstOffset,
204 VkDeviceSize size, uint32_t data)
205 {
206 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
207 cmd_buffer->device->vk.dispatch_table.CmdFillBuffer(commandBuffer, dstBuffer, dstOffset, size,
208 data);
209 }
210
211 VKAPI_ATTR void VKAPI_CALL
vkCmdDispatchIndirect(VkCommandBuffer commandBuffer,VkBuffer buffer,VkDeviceSize offset)212 vkCmdDispatchIndirect(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize offset)
213 {
214 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
215 cmd_buffer->device->vk.dispatch_table.CmdDispatchIndirect(commandBuffer, buffer, offset);
216 }
217