aobench: port to vcc

Author: Hugobros3

samples/aobench/CMakeLists.txt

@@ -1,22 +1,17 @@
-find_program(CLANG_EXE "clang")
-find_program(LLVM-SPIRV_EXE "llvm-spirv")
-
-if (CLANG_EXE AND LLVM-SPIRV_EXE)
+if (TARGET vcc)
     add_executable(aobench_host ao_host.c ao_main.c)
+    add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll COMMAND vcc ARGS ${CMAKE_CURRENT_SOURCE_DIR}/ao.comp.cpp --only-run-clang -o ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll COMMENT ao.comp.c.ll DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/ao.comp.cpp)
 
-    add_custom_command(OUTPUT ao.cl.ll COMMAND clang ARGS ${CMAKE_CURRENT_SOURCE_DIR}/ao.cl -std=clc++2021 -emit-llvm -o ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.ll -c -target spir64-unknown-unknown -O2 DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/ao.cl ${CMAKE_CURRENT_SOURCE_DIR}/ao.c)
-    add_custom_command(OUTPUT ao.cl.spv COMMAND llvm-spirv ARGS ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.ll -o ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.spv DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.ll)
-
-    set_property(SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/ao_main.c APPEND PROPERTY OBJECT_DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.spv)
+    set_property(SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/ao_main.c APPEND PROPERTY OBJECT_DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll)
     add_custom_command(TARGET aobench_host POST_BUILD
             COMMAND ${CMAKE_COMMAND} -E copy
-            ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.spv
-            ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/ao.cl.spv)
+            ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll
+            ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/ao.comp.c.ll)
 
     find_program(ISPC_EXE "ispc")
     if (ISPC_EXE)
         target_compile_definitions(aobench_host PUBLIC ENABLE_ISPC=1)
-        add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc COMMAND slim ARGS ao.cl.spv --output ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc --entry-point aobench_kernel COMMENT generating aobench.ispc DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.cl.spv)
+        add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc COMMAND slim ARGS ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll -o ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc --entry-point aobench_kernel COMMENT generating aobench.ispc DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll)
         add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc.o COMMAND ispc ARGS ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc -o ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc.o --pic -g -O2 -woff COMMENT generating aobench.ispc.o DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc)
         add_library(aobench_ispc OBJECT ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc.o)
         set_target_properties(aobench_ispc PROPERTIES LINKER_LANGUAGE C)
@@ -25,5 +20,5 @@ if (CLANG_EXE AND LLVM-SPIRV_EXE)
 
     target_link_libraries(aobench_host PRIVATE m shady runtime common)
 else()
-    message("Clang and/or llvm-spirv not found. Skipping aobench.")
+    message("Vcc not enabled. Skipping aobench.")
 endif()

samples/aobench/ao.cl

@@ -0,0 +1,58 @@
+#include <stdint.h>
+
+#define compute_shader __attribute__((annotate("shady::entry_point::Compute")))
+
+#define location(i) __attribute__((annotate("shady::location::"#i)))
+
+#define input __attribute__((address_space(389)))
+#define output __attribute__((address_space(390)))
+#define global __attribute__((address_space(1)))
+
+typedef uint32_t uvec4 __attribute__((ext_vector_type(4)));
+typedef float vec4 __attribute__((ext_vector_type(4)));
+
+typedef uint32_t uvec3 __attribute__((ext_vector_type(3)));
+typedef float vec3 __attribute__((ext_vector_type(3)));
+
+/*__attribute__((annotate("shady::builtin::FragCoord")))
+input vec4 fragCoord;
+
+location(0) input vec3 fragColor;
+location(0) output vec4 outColor;*/
+
+__attribute__((annotate("shady::builtin::WorkgroupId")))
+input uvec3 workgroup_id;
+
+__attribute__((annotate("shady::builtin::GlobalInvocationId")))
+input uvec3 global_id;
+
+float sqrtf(float) __asm__("shady::prim_op::sqrt");
+float sinf(float) __asm__("shady::prim_op::sin");
+float cosf(float) __asm__("shady::prim_op::cos");
+float fmodf(float, float) __asm__("shady::prim_op::mod");
+float fabsf(float) __asm__("shady::prim_op::abs");
+float floorf(float) __asm__("shady::prim_op::floor");
+#include "ao.c"
+
+extern "C" __attribute__((annotate("shady::workgroup_size::16::16::1")))
+compute_shader void aobench_kernel(global unsigned char* out) {
+    //outColor = (vec4) { fragColor[0], fragColor[1], fragColor[2], 1.0f };
+    //outColor = (vec4) { fragCoord[0] / 1024, fragCoord[1] / 1024, 1.0f, 1.0f };
+
+    Ctx ctx = get_init_context();
+    init_scene(&ctx);
+
+    int x = global_id.x;
+    int y = global_id.y;
+    //int x = (int) fragCoord.x % 1024;
+    //int y = (int) fragCoord.y % 1024;
+
+    // unsigned int out[3]; // = { 55, 0, 0};
+    out[0] = 255;
+    out[1] = 255;
+    render_pixel(&ctx, x + 3, y, WIDTH, HEIGHT, NSUBSAMPLES, (unsigned char*) out);
+    //out[2] = 155;
+    // out[0] = x / 4;
+    // out[1] = y / 4;
+    //outColor = (vec4) { ((int) out[0]) / 255.0f, ((int) out[1]) / 255.0f, ((int) out[2]) / 255.0f, 1.0f };
+}

samples/aobench/ao.comp.cpp

@@ -180,7 +180,7 @@ int main(int argc, char **argv) {
 #endif
 
     if (do_device || do_all) {
-        render_device(&args, img, WIDTH, HEIGHT, NSUBSAMPLES, "./ao.cl.spv");
+        render_device(&args, img, WIDTH, HEIGHT, NSUBSAMPLES, "./ao.comp.c.ll");
         saveppm("device.ppm", WIDTH, HEIGHT, img);
     }