1639{
1641 const auto& ks = *spec.
kernel;
1642
1643 bool has_image = false;
1644 for (
const auto&
b : spec.bindings) {
1645 if (
b.modality == Kakshya::DataModality::IMAGE_2D)
1646 has_image = true;
1647 }
1648
1649 std::string o;
1650 o += "#version 460\n";
1651 o += "layout(local_size_x = " + std::to_string(ws[0])
1652 + ", local_size_y = " + std::to_string(ws[1])
1653 + ", local_size_z = " + std::to_string(ws[2]) + ") in;\n\n";
1654
1655 for (
const auto&
b : spec.bindings) {
1656 if (
b.modality == Kakshya::DataModality::IMAGE_2D) {
1657 const std::string qual = (
b.direction == BindingDirection::Input)
1658 ? "readonly"
1659 : "writeonly";
1660 o +=
"layout(set = 0, binding = " + std::to_string(
b.binding_index)
1661 +
", rgba32f) " + qual +
" uniform image2D " +
b.name +
";\n";
1662 continue;
1663 }
1664 if (
b.modality == Kakshya::DataModality::TEXTURE_2D) {
1665 o +=
"layout(set = 0, binding = " + std::to_string(
b.binding_index)
1666 +
") uniform sampler2D " +
b.name +
";\n";
1667 continue;
1668 }
1669 const auto t = std::string(glsl_type(
b.format));
1670 o +=
"layout(set = 0, binding = " + std::to_string(
b.binding_index)
1671 +
", std430) buffer Block_" +
b.name
1672 +
" { " + t +
" " +
b.name +
"[]; };\n";
1673 }
1674
1676 o += "\nlayout(push_constant) uniform PC {\n";
1677 for (const auto& f : spec.pc_fields)
1678 o +=
" " +
std::string(glsl_type(f.format)) +
" " + f.
name +
";\n";
1679 o += "} pc;\n";
1680 }
1681
1682 for (const auto& f : spec.functions) {
1683 o += "\n" + f.return_type + " " + f.name + "(" + f.params + ") {\n";
1684 o += f.body;
1685 o += "\n}\n";
1686 }
1687
1688 o += "\nvoid main() {\n";
1689 o += " uint i = gl_GlobalInvocationID.x;\n";
1690 if (has_image)
1691 o += " ivec2 coord = ivec2(gl_GlobalInvocationID.xy);\n";
1692
1693 for (const auto& f : spec.pc_fields) {
1694 const auto t = std::string(glsl_type(f.format));
1695 o += " " + t + " " + f.name + " = pc." + f.name + ";\n";
1696 }
1697
1698 o += ks.body;
1699 o += "\n}\n";
1700 return o;
1701}
std::vector< PushConstantField > pc_fields
std::array< uint32_t, 3 > workgroup_size
std::optional< KernelSource > kernel
When set, KernelOp is ignored.