More complex opencl argument passing. master
authorNot Zed <notzed@gmail.com>
Sun, 29 Dec 2019 06:03:23 +0000 (16:33 +1030)
committerNot Zed <notzed@gmail.com>
Sun, 29 Dec 2019 06:03:23 +0000 (16:33 +1030)
README
opencl-args/Makefile [new file with mode: 0644]
opencl-args/compile-desc [new file with mode: 0755]
opencl-args/compute.c [new file with mode: 0644]
opencl-args/compute.h [new file with mode: 0644]
opencl-args/hello.c [new file with mode: 0644]
opencl-args/hello.cl [new file with mode: 0644]
opencl-args/zvk.c [new file with mode: 0644]
opencl-args/zvk.h [new file with mode: 0644]

diff --git a/README b/README
index ecb8cb9..c6729ef 100644 (file)
--- a/README
+++ b/README
@@ -36,6 +36,10 @@ mandelbrot-cl/
   By far most of the execution time is spent on the CPU side,
   converting the float4 array to uchar3 for output.
 
   By far most of the execution time is spent on the CPU side,
   converting the float4 array to uchar3 for output.
 
+opencl-args/
+
+  How do you pass arguments?
+
 LINKS
 -----
 
 LINKS
 -----
 
diff --git a/opencl-args/Makefile b/opencl-args/Makefile
new file mode 100644 (file)
index 0000000..bece3aa
--- /dev/null
@@ -0,0 +1,23 @@
+
+#.SUFFIXES:
+
+include ../config.make
+
+CFLAGS += -Wall -Og -g $(compute_CFLAGS)
+LDLIBS += $(compute_LDLIBS)
+
+all: hello
+
+hello.o: hello-host.h
+hello: hello.o hello-host.o compute.o zvk.o
+
+%-host.c %-host.h: %.cl.csv %.cl.spv compile-desc
+       ./compile-desc --prefix $*_ -h $*-host.h -o $*-host.c $*.cl.spv $*.cl.csv
+
+clean:
+       rm -f *.cl.csv *.cl.spv *.cl.inc
+       rm -f *-host.c *-host.h
+       rm -f *.o
+       rm -f hello
+
+.SECONDARY: hello.cl.csv hello.cl.spv
diff --git a/opencl-args/compile-desc b/opencl-args/compile-desc
new file mode 100755 (executable)
index 0000000..94bdb4c
--- /dev/null
@@ -0,0 +1,352 @@
+#!/usr/bin/perl
+
+use List::Util qw[min max];
+
+# usage: compile-desc [--prefix func-prefix] -o output.c [ -h output.h ] file.spv file.csv
+
+$prefix = "";
+$outc = "";
+$outh = "";
+@files = ();
+
+while (@ARGV) {
+    my $arg = shift @ARGV;
+    
+    if ($arg eq "--prefix") {
+       $prefix = shift @ARGV;
+    } elsif ($arg eq "-o") {
+       $outc = shift @ARGV;
+    } elsif ($arg eq "-h") {
+       $outh = shift @ARGV;
+    } else {
+       push @files, $arg;
+    }       
+}
+
+$spv = $files[0];
+$desc = $files[1];
+
+die ("Missing descriptor file") if ($desc eq "");
+die ("Missing output file") if ($outc eq "");
+
+# This takes the opencl kernel plus the descriptorsets description and ... does something with it.
+#$unit = "arguments";
+#$desc = "arguments.cl.csv";
+#$src = "arguments.cl";
+
+%info = (
+    "prefix" => $prefix,
+    );
+
+# map for primitive types which need renaming
+%type_map = (
+    "uint32", "uint32_t",
+    "int32", "int32_t",
+    );
+
+%bind_map = (
+    "pod", "ZVK_POD",
+    "buffer", "ZVK_BUFFER",
+    "ro_image", "ZVK_ROIMAGE",
+    "wo_image", "ZVK_WOIMAGE",
+    "sampler", "ZVK_SAMPLER",
+    );
+
+%kernels = ();
+
+#
+# kernel{name} = (
+#  @bindings = (
+#   index = index
+#   kind = pod | buffer
+#   sizeof = total byte size for pod
+#  )
+#  @params = (
+#   index = index
+#   name = parameter name
+#   type = opencl param type
+#   binding = binding index
+#   offset =
+#   size = parameter size for pod
+#  )
+# )
+
+# read description
+open IN,"<$desc";
+while (<IN>) {
+    chop;
+    if (m/^kernel,([^,]*),arg,([^,]*),argOrdinal,([^,]*),descriptorSet,([^,]*),binding,([^,]*),offset,([^,]*),argKind,([^,]*)(.*)/) {
+       my $kname = $1;
+       my $pname = $2;
+       my $index = $3;
+       my $set = $4;
+       my $bindid = $5;
+       my $offset = $6;
+       my $kind = $7;
+       my $rest = $8;
+       my %kern = %{$kernels{$kname}}; 
+       my %rec = %{@{$kern{params}}[$index]};
+       
+       die ("cl parser <> descriptor mismatch") if ($rec{name} != $name);
+
+       $rec{name} = $pname;
+       $rec{index} = $index;
+       $rec{binding} = $bindid;
+       $rec{offset} = $offset;
+       ($rec{size}) = ($rest =~ m/argSize,(\d*)/);
+       ($rec{type}) = ($rest =~ m/argType,([^,]*)/);
+       @{$kern{params}}[$index] = \%rec;
+       
+       my %bind = %{@{$kern{bindings}}[$bindid]};
+
+       $bind{index} = $bindid;
+       $bind{kind} = $kind;
+       if (defined $rec{size}) {
+           $bind{sizeof} = max($offset+$rec{size}, $bind{sizeof});
+       }
+                                                   
+       @{$kern{bindings}}[$bindid] = \%bind;
+
+       $kern{name} = $kname;
+       
+       $kernels{$kname} = \%kern;
+    }
+}
+close IN;
+
+# get kernel arg types, unfortunately not available in the descriptor file
+# requires strict source format:
+# kernel void\n
+# func(args,...)\n
+if (0) {
+open IN,"<$src";
+$state = 0;
+while (<IN>) {
+    if ($state == 0) {
+       if (m/^kernel void/) {
+           $state = 1;
+       }
+    } else {
+       # this only includes kernels which are already defined
+       if (m/([^\(]*)\((.*)\)/) {
+           my $kname = $1;
+           my @args = $2 =~ m/\s*([^,]+)\s*,?/g;
+           my %kern = %{$kernels{$kname}};
+           my $i = 0;
+
+           for $arg (@args) {
+               my ($type, $name) = $arg =~ m/^(.*?) ?(\w*)$/;
+               my %rec = %{$kern{params}[$i]};
+
+               print STDERR "arg: $type <- $name\n";
+               
+               $rec{type} = $type;
+
+               $kern{params}[$i] = \%rec;
+               
+               $i++;
+           }
+       }
+       $state = 0;
+    }
+}
+close IN;
+}
+
+open C,">$outc" || die("opening $outc");
+
+print C <<END;
+/* autogenerated from $desc */
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include "compute.h"
+END
+
+if ($outh ne "") {
+    open H,">$outh" || die("open $outh");
+}
+
+# Data tables - kernels
+for $kname (sort keys %kernels) {
+    my %kern = %{$kernels{$kname}};
+    my @bindings = @{$kern{bindings}};
+    my @params = @{$kern{params}};
+    
+    print C "static const struct bindinfo $info{prefix}$kern{name}_bindings[] = {\n";
+    for $binds (@bindings) {
+       my %bind = %{$binds};
+
+       print C " { $bind_map{$bind{kind}}, $bind{sizeof} },\n";
+    }
+    print C "};\n";
+
+    print C "static const struct paraminfo $info{prefix}$kern{name}_params[] = {\n";
+    for $recs (@params) {
+       my %rec = %{$recs};
+
+       print C " { $rec{binding}, $rec{offset}, $rec{size} },\n";
+    }
+    print C "};\n";
+
+    my $nbindings = $#bindings + 1;
+    my $nparams = $#params + 1;
+
+    print C <<END;
+static const struct kerninfo $info{prefix}$kern{name} = {
+  .name = "$kern{name}",
+  .nbindings = $nbindings,
+  .nparams = $nparams,
+  .bindings = $info{prefix}$kern{name}_bindings,
+  .params = $info{prefix}$kern{name}_params
+ };
+END
+}
+
+# Data tables - module
+open SPV,"<:raw", $spv;
+my $i = 0;
+print C "/* $spv */\n";
+print C "static const uint32_t module_spv[] = {\n";
+while (read SPV,my $code,4) {
+    printf C " 0x%08x,", unpack("V*", $code);
+    if (($i++) % 8 == 7) {
+       print C "\n";
+    }
+}
+print C "\n};\n";
+close SPV;
+
+#
+# module definition
+#
+$nkernels = keys %kernels;
+print C <<END;
+const struct modinfo $info{prefix}modinfo = {
+ .code = module_spv,
+ .codeSize = sizeof(module_spv),
+ .nkernels = $nkernels,
+END
+if (H) {
+    print H "extern const struct modinfo $info{prefix}modinfo;\n";
+}
+     
+print C " .kernels = {\n";
+for $kname (sort keys %kernels) {
+    my %kern = %{$kernels{$kname}};
+
+    print C "  &$info{prefix}$kern{name},\n";
+}
+print C " },\n";
+print C "};\n";
+
+#
+# create method to instantiate 1 kernel
+#
+if (0) {
+    for $kname (sort keys %kernels) {
+       my %kern = %{$kernels{$kname}};
+       my @bindings = @{$kern{bindings}};
+       my @params = @{$kern{params}};
+       
+       my @buffers;
+       my @pods;
+       my @images;
+    
+       print C "struct kernstate *$info{prefix}$kern{name}_create(struct zvk *zvk) {\n";
+       print C " return compute_create(zvk, &$info{prefix}$kern{name});\n";
+       print C "}\n\n";
+
+       if (H) {
+           print H "struct kernstate *$info{prefix}$kern{name}_create(struct zvk *zvk);\n";
+       }
+    }
+}
+
+# create method to update arguments
+#
+for $kname (sort keys %kernels) {
+    my %kern = %{$kernels{$kname}};
+    my @bindings = @{$kern{bindings}};
+    my @params = @{$kern{params}};
+
+    my @buffers;
+    my @pods;
+    my @images;
+
+    my $proto = "void $info{prefix}$kern{name}_setArgs(struct kernstate *state";
+        
+    for $recs (@params) {
+       my %rec = %{$recs};
+       my %bind = %{$bindings[$rec{binding}]};
+
+       if ($bind{kind} eq "buffer") {
+           $proto .= ", VkDescriptorBufferInfo $rec{name}";
+           push @buffers, \%rec;
+       } elsif ($bind{kind} eq "pod") {
+           if (defined($type_map{$rec{type}})) {
+               $proto .= ", $type_map{$rec{type}} $rec{name}";
+               push @pods, \%rec;
+           } else {
+               $proto .= ", $rec{type} $rec{name}";
+               push @pods, \%rec;
+           }
+       } elsif ($bind{kind} eq "ro_image" || $bind{kind} eq "wo_image") {
+           $proto .= ", VkDescriptorImageInfo $rec{name}";
+           push @images, \%rec;
+       }       
+    }
+    $proto .= ")";
+    print C "$proto {\n";
+    if (H) {
+       print H "$proto;\n";
+    }
+    print C " struct zvk *zvk = state->mod->zvk;\n";
+
+    if (@buffers) {
+       print C " /* update buffer pointers */\n";
+       print C " VkDescriptorBufferInfo *bufferSet[] = {\n";
+       for $recs (@buffers) {
+           my %rec = %{$recs};
+           print C "  &$rec{name},\n";
+       }
+       print C " };\n\n";
+       print C " compute_setBuffers(state, bufferSet);\n";
+    }
+
+    if (@images) {
+       print C " /* update image pointers */\n";
+       print C " VkDescriptorImageInfo *imageSet[] = {\n";
+       for $recs (@images) {
+           my %rec = %{$recs};
+           print C "  &$rec{name},\n";
+       }
+       print C " };\n";
+       print C " compute_setImages(state, imageSet);\n";
+    }
+
+    if (@pods) {
+       my @bindpods = grep { %{$_}{kind} eq "pod" } @bindings;
+       my $podSize = %{$bindpods[0]}{sizeof};
+       
+       print C " /* update pod memory */\n";
+       print C " void *pod __attribute__ ((aligned(16)));\n";
+       print C " ZVK_FATAL(vkMapMemory(zvk->device, state->podMemory, 0, $podSize, 0, &pod));\n";
+       
+       for $recs (@pods) {
+           my %rec = %{$recs};
+
+           print C " memcpy(pod + $rec{offset}, &$rec{name}, $rec{size});\n";
+       }
+
+       print C " vkUnmapMemory(zvk->device, state->podMemory);\n";
+    }
+    
+    print C "}\n\n";
+}
+
+if (H) {
+    close H;
+}
+close C;
+
diff --git a/opencl-args/compute.c b/opencl-args/compute.c
new file mode 100644 (file)
index 0000000..61af5ee
--- /dev/null
@@ -0,0 +1,314 @@
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include "compute.h"
+
+/* ********************************************************************** */
+
+static int mapDescriptorType(int binding) {
+       switch (binding) {
+       case ZVK_BUFFER:
+       case ZVK_POD:
+               return VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
+       case ZVK_ROIMAGE:
+               return VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE;
+       case ZVK_WOIMAGE:
+               return VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
+       case ZVK_SAMPLER:
+               return VK_DESCRIPTOR_TYPE_SAMPLER;
+       }
+       return -1;
+}
+
+struct modstate *compute_createModule(struct zvk *zvk, const struct modinfo *mod) {
+       struct modstate *ms = calloc(1, sizeof(*ms));
+
+       ms->zvk = zvk;
+       ms->modinfo = mod;
+
+       VkShaderModuleCreateInfo createInfo = {
+               .sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO,
+               .codeSize = mod->codeSize,
+               .pCode = mod->code,
+       };
+
+       ZVK_FATAL(vkCreateShaderModule(zvk->device, &createInfo, NULL, &ms->shader));
+
+       return ms;
+}
+
+void compute_destroyModule(struct modstate *mod) {
+       vkDestroyShaderModule(mod->zvk->device, mod->shader, NULL);
+       free(mod);
+}
+
+static int kernelIndex(const struct modinfo *mi, const char * name) {
+       for (int i=0;i<mi->nkernels;i++) {
+               if (strcmp(mi->kernels[i]->name, name) == 0)
+                       return i;
+       }
+       return -1;
+}
+
+struct kernstate *compute_createKernel(struct modstate *mod, const char *name) {
+       int ki = kernelIndex(mod->modinfo, name);
+
+       if (ki < 0)
+               return NULL;
+
+       struct zvk *zvk = mod->zvk;
+       struct kernstate *ks = calloc(1, sizeof(*ks));
+       const struct kerninfo *kern = mod->modinfo->kernels[ki];
+
+       ks->mod = mod;
+       ks->kern = kern;
+
+       /* allocate data bindings */
+       VkDescriptorSetLayoutBinding layout_bindings[kern->nbindings];
+       VkDescriptorBindingFlagBitsEXT layout_bindings_flags[kern->nbindings];
+
+       for (int i=0;i<kern->nbindings;i++) {
+               memset(&layout_bindings[i], 0, sizeof(layout_bindings[0]));
+               layout_bindings_flags[i] = VK_DESCRIPTOR_BINDING_UPDATE_AFTER_BIND_BIT_EXT;
+               layout_bindings[i].binding = i;
+               layout_bindings[i].descriptorType = mapDescriptorType(kern->bindings[i].type);
+               layout_bindings[i].descriptorCount = 1;
+               layout_bindings[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+       }
+       VkDescriptorSetLayoutBindingFlagsCreateInfoEXT flags = {
+               .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_BINDING_FLAGS_CREATE_INFO_EXT,
+               .bindingCount = kern->nbindings,
+               .pBindingFlags = layout_bindings_flags,
+       };
+       VkDescriptorSetLayoutCreateInfo descriptor_layout = {
+               .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
+               .pNext = &flags,
+               .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_UPDATE_AFTER_BIND_POOL_BIT_EXT,
+               .bindingCount = kern->nbindings,
+               .pBindings = layout_bindings,
+       };
+
+       ZVK_FATAL(vkCreateDescriptorSetLayout(zvk->device, &descriptor_layout, NULL, &ks->descriptorSetLayout));
+
+       VkDescriptorSetAllocateInfo alloc_info[] = {
+               {
+                       .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO,
+                       .descriptorPool = zvk->descriptorPool,
+                       .descriptorSetCount = 1,
+                       .pSetLayouts = &ks->descriptorSetLayout,
+               },
+       };
+
+       ZVK_FATAL(vkAllocateDescriptorSets(zvk->device, alloc_info, ks->descriptorSets));
+
+       /* Check for the (1) 'pod' binding, allocate memory for it */
+       for (int i=0;i<kern->nbindings;i++) {
+               if (kern->bindings[i].type == ZVK_POD) {
+                       zvkAllocBuffer(zvk, kern->bindings[i].size,
+                                      VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                                      VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
+                                      &ks->podBuffer, &ks->podMemory);
+
+                       VkDescriptorBufferInfo bufferInfo = {
+                               .buffer = ks->podBuffer,
+                               .offset = 0,
+                               .range = kern->bindings[i].size,
+                       };
+                       VkWriteDescriptorSet writeSet = {
+                               .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+                               .dstSet = ks->descriptorSets[0],
+                               .dstBinding = i,
+                               .descriptorCount = 1,
+                               .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
+                               .pBufferInfo = &bufferInfo,
+                       };
+                       vkUpdateDescriptorSets(zvk->device, 1, &writeSet, 0, NULL);
+                       break;
+               }
+       }
+
+       /* Create pipeline */
+       VkPipelineLayoutCreateInfo pipelineinfo = {
+               .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
+               .setLayoutCount = 1,
+               .pSetLayouts = &ks->descriptorSetLayout,
+       };
+
+       ZVK_FATAL(vkCreatePipelineLayout(zvk->device, &pipelineinfo, NULL, &ks->pipelineLayout));
+
+       VkComputePipelineCreateInfo pipeline = {
+               .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
+               .stage =
+               {
+                       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
+                       .stage = VK_SHADER_STAGE_COMPUTE_BIT,
+                       .module = mod->shader,
+                       .pName = kern->name,
+               },
+               .layout = ks->pipelineLayout
+
+       };
+       ZVK_FATAL(vkCreateComputePipelines(zvk->device, NULL, 1, &pipeline, NULL, &ks->pipeline));
+
+       return ks;
+}
+
+void compute_destroyKernel(struct kernstate *ks) {
+       struct zvk *zvk = ks->mod->zvk;
+
+       if (ks->podMemory) {
+               vkFreeMemory(zvk->device, ks->podMemory, NULL);
+               vkDestroyBuffer(zvk->device, ks->podBuffer, NULL);
+       }
+
+       vkDestroyPipeline(zvk->device, ks->pipeline, NULL);
+       vkDestroyPipelineLayout(zvk->device, ks->pipelineLayout, NULL);
+
+       vkFreeDescriptorSets(zvk->device, zvk->descriptorPool, 1, ks->descriptorSets);
+       vkDestroyDescriptorSetLayout(zvk->device, ks->descriptorSetLayout, NULL);
+
+       free(ks);
+}
+
+// one time optional?
+VkCommandBuffer compute_createCommand(struct kernstate *ks, uint32_t sizex, uint32_t sizey, uint32_t sizez) {
+       struct zvk *zvk = ks->mod->zvk;
+       VkCommandBuffer commandBuffers[1];
+
+       /* Create a command buffer to run this kernel with it's data set for the given size */
+       VkCommandBufferAllocateInfo cmdinfo = {
+               .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO,
+               .commandPool = zvk->commandPool,
+               .level = VK_COMMAND_BUFFER_LEVEL_PRIMARY,
+               .commandBufferCount = 1,
+       };
+
+       ZVK_FATAL(vkAllocateCommandBuffers(zvk->device, &cmdinfo, commandBuffers));
+
+       VkCommandBufferBeginInfo beginInfo = {
+               .sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
+               //.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT,
+               .flags = 0,
+       };
+       ZVK_FATAL(vkBeginCommandBuffer(commandBuffers[0], &beginInfo));
+
+       vkCmdBindPipeline(commandBuffers[0], VK_PIPELINE_BIND_POINT_COMPUTE, ks->pipeline);
+       vkCmdBindDescriptorSets(commandBuffers[0], VK_PIPELINE_BIND_POINT_COMPUTE, ks->pipelineLayout, 0, 1, ks->descriptorSets, 0, NULL);
+
+       vkCmdDispatch(commandBuffers[0], sizex, sizey, sizez);
+
+       ZVK_FATAL(vkEndCommandBuffer(commandBuffers[0]));
+
+       return commandBuffers[0];
+}
+
+// rather inefficient way to set arguments one at a time
+void compute_setArg(struct kernstate *ks, int index, void *data, size_t size) {
+       struct zvk *zvk = ks->mod->zvk;
+       const struct kerninfo *kern = ks->kern;
+       const struct paraminfo *pi = &kern->params[index];
+       const struct bindinfo *bi = &kern->bindings[pi->binding];
+
+       switch (bi->type) {
+       case ZVK_POD: {
+               void *pod __attribute__ ((aligned(16)));
+
+               ZVK_FATAL(!(size == pi->size));
+               ZVK_FATAL(vkMapMemory(zvk->device, ks->podMemory, 0, VK_WHOLE_SIZE, 0, &pod));
+               memcpy(pod + pi->offset, data, size);
+               vkUnmapMemory(zvk->device, ks->podMemory);
+               break;
+       }
+       case ZVK_BUFFER:
+               ZVK_FATAL(!(size == sizeof(VkDescriptorBufferInfo)));
+               VkWriteDescriptorSet bufferSet = {
+                       .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+                       .dstSet = ks->descriptorSets[0],
+                       .dstBinding = pi->binding,
+                       .descriptorCount = 1,
+                       .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
+                       .pBufferInfo = data
+               };
+               vkUpdateDescriptorSets(zvk->device, 1, &bufferSet, 0, NULL);
+               break;
+       case ZVK_WOIMAGE:
+               ZVK_FATAL(!(size == sizeof(VkDescriptorImageInfo)));
+               VkWriteDescriptorSet woSet = {
+                       .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+                       .dstSet = ks->descriptorSets[0],
+                       .dstBinding = pi->binding,
+                       .descriptorCount = 1,
+                       .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
+                       .pImageInfo = data
+               };
+               vkUpdateDescriptorSets(zvk->device, 1, &woSet, 0, NULL);
+               break;
+       case ZVK_ROIMAGE:
+               ZVK_FATAL(!(size == sizeof(VkDescriptorImageInfo)));
+               VkWriteDescriptorSet roSet = {
+                       .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+                       .dstSet = ks->descriptorSets[0],
+                       .dstBinding = pi->binding,
+                       .descriptorCount = 1,
+                       .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
+                       .pImageInfo = data
+               };
+               vkUpdateDescriptorSets(zvk->device, 1, &roSet, 0, NULL);
+               break;
+               // shader?
+       }
+}
+
+void compute_setBuffers(struct kernstate *ks, VkDescriptorBufferInfo *buffers[]) {
+       VkWriteDescriptorSet writeSet = {
+               .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+               .dstSet = ks->descriptorSets[0],
+               .descriptorCount = 1,
+               .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
+       };
+       const struct kerninfo *kern = ks->kern;
+       struct zvk *zvk = ks->mod->zvk;
+
+       for (int i=0,bindex=0;i<kern->nparams;i++) {
+               const struct paraminfo *pi = &kern->params[i];
+               const struct bindinfo *bi = &kern->bindings[pi->binding];
+
+               if (bi->type == ZVK_BUFFER) {
+                       writeSet.pBufferInfo = buffers[bindex];
+                       writeSet.dstBinding = pi->binding;
+
+                       vkUpdateDescriptorSets(zvk->device, 1, &writeSet, 0, NULL);
+                       bindex++;
+               }
+       }
+}
+
+void compute_setImages(struct kernstate *ks, VkDescriptorImageInfo *images[]) {
+       VkWriteDescriptorSet writeSet = {
+               .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
+               .dstSet = ks->descriptorSets[0],
+               .descriptorCount = 1,
+       };
+       const struct kerninfo *kern = ks->kern;
+       struct zvk *zvk = ks->mod->zvk;
+
+       for (int i=0,bindex=0;i<kern->nparams;i++) {
+               const struct paraminfo *pi = &kern->params[i];
+               const struct bindinfo *bi = &kern->bindings[pi->binding];
+
+               if (bi->type == ZVK_WOIMAGE) {
+                       writeSet.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
+               } else if (bi->type == ZVK_ROIMAGE) {
+                       writeSet.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE;
+               } else
+                       continue;
+
+               writeSet.pImageInfo = images[bindex];
+               writeSet.dstBinding = pi->binding;
+
+               vkUpdateDescriptorSets(zvk->device, 1, &writeSet, 0, NULL);
+               bindex++;
+       }
+}
diff --git a/opencl-args/compute.h b/opencl-args/compute.h
new file mode 100644 (file)
index 0000000..203e34f
--- /dev/null
@@ -0,0 +1,71 @@
+
+#ifndef COMPUTE_H
+#define COMPUTE_H
+
+#include "zvk.h"
+
+/* ********************************************************************** */
+
+struct paraminfo {
+       int binding;
+       int offset;
+       int size;
+};
+
+struct bindinfo {
+       int type;
+       int size;
+};
+
+struct kerninfo {
+       const char *name;
+       int nbindings;
+       int nparams;
+       const struct bindinfo *bindings;
+       const struct paraminfo *params;
+};
+
+struct kernstate {
+       struct modstate *mod;
+       const struct kerninfo *kern;
+
+       VkDeviceMemory podMemory;
+       VkBuffer podBuffer;
+
+       VkDescriptorSetLayout descriptorSetLayout;
+       VkDescriptorSet descriptorSets[1];
+
+       VkPipelineLayout pipelineLayout;
+       VkPipeline pipeline;
+};
+
+/* ********************************************************************** */
+
+struct modinfo {
+       const uint32_t *code;
+       size_t codeSize;
+       int nkernels;
+       const struct kerninfo *kernels[];
+};
+
+struct modstate {
+       struct zvk *zvk;
+       const struct modinfo *modinfo;
+
+       VkShaderModule shader;
+};
+
+struct modstate *compute_createModule(struct zvk *zvk, const struct modinfo *mod);
+void compute_destroyModule(struct modstate *mod);
+
+struct kernstate *compute_createKernel(struct modstate *mod, const char *name);
+void compute_destroyKernel(struct kernstate *ks);
+
+VkCommandBuffer compute_createCommand(struct kernstate *ks, uint32_t sizex, uint32_t sizey, uint32_t sizez);
+
+void compute_setArg(struct kernstate *ks, int index, void *data, size_t size);
+
+void compute_setBuffers(struct kernstate *ks, VkDescriptorBufferInfo *buffers[]);
+void compute_setImages(struct kernstate *ks, VkDescriptorImageInfo *images[]);
+
+#endif
diff --git a/opencl-args/hello.c b/opencl-args/hello.c
new file mode 100644 (file)
index 0000000..0c5a723
--- /dev/null
@@ -0,0 +1,125 @@
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#include "compute.h"
+
+#include "hello-host.h"
+
+void zvkAllocBuffer2(struct zvk *g, VkDeviceSize dataSize, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, struct zvkBuffer *buffer) {
+       zvkAllocBuffer(g, dataSize, usage, properties, &buffer->info.buffer, &buffer->memory);
+       buffer->info.offset = 0;
+       buffer->info.range = dataSize;
+}
+
+void zvkFreeBuffer(struct zvk *zvk, struct zvkBuffer *buffer) {
+       vkFreeMemory(zvk->device, buffer->memory, NULL);
+       vkDestroyBuffer(zvk->device, buffer->info.buffer, NULL);
+}
+
+void zvkSetBuffer(struct zvk *zvk, struct zvkBuffer *buffer, void *data, size_t size) {
+       void *mem;
+
+       vkMapMemory(zvk->device, buffer->memory, 0, buffer->info.range, 0, &mem);
+       memcpy(mem, data, size);
+       vkUnmapMemory(zvk->device, buffer->memory);
+}
+
+int main(int argc, char **argv) {
+       struct zvk *zvk = zvkInit();
+       struct modstate *mod = compute_createModule(zvk, &hello_modinfo);
+       struct kernstate *hello = compute_createKernel(mod, "hello");
+       struct zvkBuffer dstBuffer, srcBuffer;
+
+       zvkAllocBuffer2(zvk, 64, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                       VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
+                       &dstBuffer);
+
+       zvkAllocBuffer2(zvk, 13, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                       VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
+                       &srcBuffer);
+
+       zvkSetBuffer(zvk, &srcBuffer, "Hello World!\n", 13);
+
+       hello_hello_setArgs(hello, dstBuffer.info, 64, srcBuffer.info, 13);
+
+       VkCommandBuffer cmd = compute_createCommand(hello, 64, 1, 1);
+
+       {
+               VkResult res;
+               VkSubmitInfo submitInfo[] = {
+                       {
+                               .sType = VK_STRUCTURE_TYPE_SUBMIT_INFO,
+                               .commandBufferCount = 1,
+                               .pCommandBuffers = &cmd,
+                       },
+               };
+
+               VkFence fence;
+               VkFenceCreateInfo fenceInfo = {
+                       .sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO
+               };
+               ZVK_FATAL(vkCreateFence(zvk->device, &fenceInfo, NULL, &fence));
+
+               ZVK_FATAL(vkQueueSubmit(zvk->computeQueue, ZVK_ARRAY_SIZEOF(submitInfo), submitInfo, fence));
+
+               do {
+                       res = vkWaitForFences(zvk->device, 1, &fence, VK_TRUE, 1000000);
+               } while (res == VK_TIMEOUT);
+
+               vkDestroyFence(zvk->device, fence, NULL);
+
+               char *dst;
+               printf("dst:\n");
+               vkMapMemory(zvk->device, dstBuffer.memory, 0, dstBuffer.info.range, 0, (void **)&dst);
+               fwrite(dst, 1, dstBuffer.info.range, stdout);
+               printf("\n");
+               vkUnmapMemory(zvk->device, dstBuffer.memory);
+       }
+
+       hello_hello_setArgs(hello, dstBuffer.info, 64, srcBuffer.info, 6);
+
+
+       {
+               VkResult res;
+               VkSubmitInfo submitInfo[] = {
+                       {
+                               .sType = VK_STRUCTURE_TYPE_SUBMIT_INFO,
+                               .commandBufferCount = 1,
+                               .pCommandBuffers = &cmd,
+                       },
+               };
+
+               VkFence fence;
+               VkFenceCreateInfo fenceInfo = {
+                       .sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO
+               };
+               ZVK_FATAL(vkCreateFence(zvk->device, &fenceInfo, NULL, &fence));
+
+               ZVK_FATAL(vkQueueSubmit(zvk->computeQueue, ZVK_ARRAY_SIZEOF(submitInfo), submitInfo, fence));
+
+               do {
+                       res = vkWaitForFences(zvk->device, 1, &fence, VK_TRUE, 1000000);
+               } while (res == VK_TIMEOUT);
+
+               vkDestroyFence(zvk->device, fence, NULL);
+
+               char *dst;
+               printf("dst:\n");
+               vkMapMemory(zvk->device, dstBuffer.memory, 0, dstBuffer.info.range, 0, (void **)&dst);
+               fwrite(dst, 1, dstBuffer.info.range, stdout);
+               printf("\n");
+               vkUnmapMemory(zvk->device, dstBuffer.memory);
+       }
+
+
+       vkFreeCommandBuffers(zvk->device, zvk->commandPool, 1, &cmd);
+
+       zvkFreeBuffer(zvk, &srcBuffer);
+       zvkFreeBuffer(zvk, &dstBuffer);
+
+       compute_destroyKernel(hello);
+       compute_destroyModule(mod);
+       zvkDestroy(zvk);
+}
diff --git a/opencl-args/hello.cl b/opencl-args/hello.cl
new file mode 100644 (file)
index 0000000..b04f099
--- /dev/null
@@ -0,0 +1,8 @@
+
+kernel void hello(global char *dst, uint dsize, constant char *src, uint ssize) {
+       uint x = get_global_id(0);
+
+       if (x < dsize) {
+               dst[x] = src[x % ssize];
+       }
+}
diff --git a/opencl-args/zvk.c b/opencl-args/zvk.c
new file mode 100644 (file)
index 0000000..04977ce
--- /dev/null
@@ -0,0 +1,282 @@
+ /*
+The MIT License (MIT)
+
+Copyright (C) 2017 Eric Arnebäck
+Copyright (C) 2019 Michael Zucchi
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+
+ */
+
+/*
+  Convert a few re-usable parts of compute-only code to a library.
+
+  It removes a little bit of the clutter in later experiments.
+ */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+#include "zvk.h"
+
+/**
+ * This finds the memory type index for the memory on a specific device.
+ */
+static int find_memory_type(VkPhysicalDeviceMemoryProperties *memory, uint32_t typeMask, VkMemoryPropertyFlagBits query)
+{
+       for (int i = 0; i < memory->memoryTypeCount; i++) {
+               if (((1 << i) & typeMask) && (memory->memoryTypes[i].propertyFlags & query) == query)
+                       return i;
+       }
+       return -1;
+}
+
+/**
+ * Create vulkan instance.
+ * 
+ * Compute-only requires no extensions so this is simple.
+ * 
+ * To turn on debugging layers:
+ *  export VK_INSTANCE_LAYERS=VK_LAYER_LUNARG_standard_validation
+ */
+static void init_instance(struct zvk *g)
+{
+       VkApplicationInfo app = {
+               .sType = VK_STRUCTURE_TYPE_APPLICATION_INFO,
+               .pApplicationName = "mandelbrot",
+               .pEngineName = "none",
+               .apiVersion = VK_API_VERSION_1_0,
+       };
+       const char * const extensions[] = {
+               "VK_KHR_get_physical_device_properties2"
+       };
+       VkInstanceCreateInfo info = {
+               .sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO,
+               .pApplicationInfo = &app,
+               .enabledExtensionCount = sizeof(extensions) / sizeof(*extensions),
+               .ppEnabledExtensionNames = extensions,
+       };
+
+       ZVK_FATAL(vkCreateInstance(&info, NULL, &g->instance));
+}
+
+/**
+ * This finds a suitable device and queue family.
+ * 
+ * In this case it is a device that supports a compute queue.  It
+ * preferentially looks for a non-graphics compute queue.
+ * 
+ * It could also make use of:
+ *  VkPhysicalDeviceProperties props;
+ *  vkGetPhysicalDeviceProperties(devs[i], &props);
+ */
+static void init_device(struct zvk *g)
+{
+       uint32_t devcount;
+
+       vkEnumeratePhysicalDevices(g->instance, &devcount, NULL);
+       VkPhysicalDevice devs[devcount];
+       ZVK_FATAL(vkEnumeratePhysicalDevices(g->instance, &devcount, devs));
+
+       int best = 0;
+       int devid = -1;
+       int queueid = -1;
+
+       for (int i = 0; i < devcount; i++) {
+               uint32_t family_count;
+
+               vkGetPhysicalDeviceQueueFamilyProperties(devs[i], &family_count, NULL);
+               VkQueueFamilyProperties famprops[family_count];
+               vkGetPhysicalDeviceQueueFamilyProperties(devs[i], &family_count, famprops);
+
+               for (uint32_t j = 0; j < family_count; j++) {
+                       int score = 0;
+
+                       if ((famprops[j].queueFlags & VK_QUEUE_COMPUTE_BIT) != 0)
+                               score += 1;
+                       if ((famprops[j].queueFlags & VK_QUEUE_GRAPHICS_BIT) == 0)
+                               score += 1;
+
+                       if (score > best) {
+                               score = best;
+                               devid = i;
+                               queueid = j;
+                       }
+               }
+       }
+
+       if (devid == -1) {
+               fprintf(stderr, "Unable to find suitable device\n");
+               exit(1);
+       }
+
+       g->physicalDevice = devs[devid];
+       g->computeDeviceIndex = devid;
+       g->computeQueueIndex = queueid;
+#if 0
+       {
+               VkExtensionProperties prop[200];
+               uint32_t count = 200;
+               vkEnumerateDeviceExtensionProperties(g->physicalDevice, NULL, &count, prop);
+               for (int i=0;i<count;i++) {
+                       printf("  %s\n", prop[i].extensionName);
+               }
+               fflush(stdout);
+       }
+       {
+               VkPhysicalDeviceShaderFloat16Int8FeaturesKHR ff = {
+                       .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_FLOAT16_INT8_FEATURES_KHR,
+               };
+               VkPhysicalDeviceFeatures2 feat = {
+                       .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2,
+                       .pNext = &ff,
+               };
+               vkGetPhysicalDeviceFeatures2(g->physicalDevice, &feat);
+
+               printf("shaderFloat16 = %d\n", ff.shaderFloat16);
+               printf("shaderInt8 = %d\n", ff.shaderInt8);
+       }
+       {
+               VkPhysicalDeviceProperties2 prop = {
+                       .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2,
+               };
+               vkGetPhysicalDeviceProperties2(g->physicalDevice, &prop);
+
+       }
+#endif
+       float qpri[] = {0.0f};
+       VkDeviceQueueCreateInfo qinfo = {
+               .sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO,
+               .queueCount = 1,
+               .pQueuePriorities = qpri, // Note: cannot be null
+               .queueFamilyIndex = g->computeQueueIndex,
+       };
+       const char * const extensions[] = {
+               VK_KHR_SHADER_FLOAT16_INT8_EXTENSION_NAME,
+               "VK_EXT_descriptor_indexing",
+               "VK_KHR_maintenance3",
+       };
+       VkPhysicalDeviceDescriptorIndexingFeaturesEXT di = {
+               .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_DESCRIPTOR_INDEXING_FEATURES_EXT,
+               .descriptorBindingStorageBufferUpdateAfterBind = 1,
+       };
+       VkPhysicalDeviceShaderFloat16Int8FeaturesKHR ff = {
+               .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_FLOAT16_INT8_FEATURES_KHR,
+               .pNext = &di,
+               .shaderInt8 = 1,
+       };
+       VkDeviceCreateInfo devinfo = {
+               .sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO,
+               .pNext = &ff,
+               .queueCreateInfoCount = 1,
+               .pQueueCreateInfos = &qinfo,
+               .enabledExtensionCount = ZVK_ARRAY_SIZEOF(extensions),
+               .ppEnabledExtensionNames = extensions,
+       };
+
+       ZVK_FATAL(vkCreateDevice(g->physicalDevice, &devinfo, NULL, &g->device));
+
+       /* These values are cached for convenience */
+       vkGetPhysicalDeviceMemoryProperties(g->physicalDevice, &g->deviceMemoryProperties);
+       vkGetDeviceQueue(g->device, g->computeQueueIndex, 0, &g->computeQueue);
+}
+
+/**
+ * Create a descriptor pool for later use.
+ */
+static void init_descriptor_pool(struct zvk *g)
+{
+       VkDescriptorPoolSize type_count[] = {
+               {
+                       .type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
+                       .descriptorCount = 32,
+               }
+       };
+       VkDescriptorPoolCreateInfo descriptor_pool = {
+               .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO,
+               .flags = VK_DESCRIPTOR_POOL_CREATE_FREE_DESCRIPTOR_SET_BIT
+                       | VK_DESCRIPTOR_POOL_CREATE_UPDATE_AFTER_BIND_BIT_EXT,
+               .maxSets = 1,
+               .poolSizeCount = ZVK_ARRAY_SIZEOF(type_count),
+               .pPoolSizes = type_count,
+       };
+
+       ZVK_FATAL(vkCreateDescriptorPool(g->device, &descriptor_pool, NULL, &g->descriptorPool));
+}
+
+/**
+ * Create a command pool for later use.
+ */
+static void init_command_pool(struct zvk *g)
+{
+       VkCommandPoolCreateInfo poolinfo = {
+               .sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO,
+               .queueFamilyIndex = g->computeQueueIndex,
+       };
+
+       ZVK_FATAL(vkCreateCommandPool(g->device, &poolinfo, NULL, &g->commandPool));
+}
+
+struct zvk *zvkInit(void) {
+       struct zvk *state = calloc(1, sizeof(*state));
+
+       init_instance(state);
+       init_device(state);
+       init_descriptor_pool(state);
+       init_command_pool(state);
+
+       return state;
+}
+
+void zvkDestroy(struct zvk *zvk) {
+       vkDestroyCommandPool(zvk->device, zvk->commandPool, NULL);
+       vkDestroyDescriptorPool(zvk->device, zvk->descriptorPool, NULL);
+       vkDestroyDevice(zvk->device, NULL);
+       vkDestroyInstance(zvk->instance, NULL);
+       free(zvk);
+}
+
+/**
+ * Buffers are created in three steps:
+ * 1) create buffer, specifying usage and size
+ * 2) allocate memory based on memory requirements
+ * 3) bind memory
+ *
+ */
+void zvkAllocBuffer(struct zvk *g, VkDeviceSize dataSize, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer *buffer, VkDeviceMemory *memory){
+       VkMemoryRequirements req;
+       VkBufferCreateInfo buf_info = {
+               .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+               .size = dataSize,
+               .usage = usage,
+               .sharingMode = VK_SHARING_MODE_EXCLUSIVE
+       };
+       ZVK_FATAL(vkCreateBuffer(g->device, &buf_info, NULL, buffer));
+
+       vkGetBufferMemoryRequirements(g->device, *buffer, &req);
+
+       VkMemoryAllocateInfo alloc = {
+               .sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO,
+               .allocationSize = req.size,
+               .memoryTypeIndex = find_memory_type(&g->deviceMemoryProperties, req.memoryTypeBits, properties)
+       };
+       ZVK_FATAL(vkAllocateMemory(g->device, &alloc, NULL, memory));
+
+       ZVK_FATAL(vkBindBufferMemory(g->device, *buffer, *memory, 0));
+}
diff --git a/opencl-args/zvk.h b/opencl-args/zvk.h
new file mode 100644 (file)
index 0000000..70afd44
--- /dev/null
@@ -0,0 +1,52 @@
+
+#ifndef ZVK_H
+#define ZVK_H
+
+#include <vulkan/vulkan.h>
+//#include <stdlib.h>
+//#include <stdio.h>
+
+struct zvk {
+       VkInstance instance;
+       VkPhysicalDevice physicalDevice;
+       VkDevice device;
+       VkPhysicalDeviceMemoryProperties deviceMemoryProperties;
+
+       VkDescriptorPool descriptorPool;
+       VkCommandPool commandPool;
+
+       uint32_t computeQueueIndex;
+       uint32_t computeDeviceIndex;
+
+       VkQueue computeQueue;
+};
+
+struct zvkBuffer {
+       VkDescriptorBufferInfo info;
+       VkDeviceMemory memory;
+};
+
+#define ZVK_BUFFER 0
+#define ZVK_POD 1
+#define ZVK_ROIMAGE 2
+#define ZVK_WOIMAGE 3
+#define ZVK_SAMPLER 4
+
+#define ZVK_FATAL(r) do {                                      \
+       VkResult res = (r);                                     \
+        if (res) {                                             \
+               fprintf(stderr, "failed: %s = %d", #r, res);    \
+               abort();                                        \
+       }                                                       \
+} while (0)
+
+struct zvk *zvkInit(void);
+void zvkAllocBuffer(struct zvk *g, VkDeviceSize dataSize, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer *buffer, VkDeviceMemory *memory);
+void zvkAllocBuffer2(struct zvk *g, VkDeviceSize dataSize, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, struct zvkBuffer *buffer);
+
+void zvkDestroy(struct zvk *zvk);
+
+
+#define ZVK_ARRAY_SIZEOF(a) (sizeof(a)/sizeof(*a))
+
+#endif