added flag to set word and ptr size

Hugobros3 · Hugobros3 · commit 50aa70ee523d · 2024-04-28T11:41:47.000+02:00
diff --git a/include/shady/ir.h b/include/shady/ir.h
@@ -64,6 +64,18 @@ const Node* get_declaration(const Module*, String);
 
 //////////////////////////////// IR Arena ////////////////////////////////
 
+typedef struct {
+    IntSizes ptr_size;
+    /// The base type for emulated memory
+    IntSizes word_size;
+} PointerModel;
+
+typedef struct {
+    PointerModel memory;
+} TargetConfig;
+
+TargetConfig default_target_config();
+
 typedef struct {
     bool name_bound;
     bool check_op_classes;
@@ -89,11 +101,7 @@ typedef struct {
         uint32_t workgroup_size[3];
     } specializations;
 
-    struct {
-        IntSizes ptr_size;
-        /// The base type for emulated memory
-        IntSizes word_size;
-    } memory;
+    PointerModel memory;
 
     /// 'folding' optimisations - happen in the constructors directly
     struct {
@@ -102,7 +110,7 @@ typedef struct {
     } optimisations;
 } ArenaConfig;
 
-ArenaConfig default_arena_config();
+ArenaConfig default_arena_config(const TargetConfig* target);
 
 IrArena* new_ir_arena(ArenaConfig);
 void destroy_ir_arena(IrArena*);
@@ -314,6 +322,8 @@ typedef struct CompilerConfig_ {
         uint32_t subgroup_size;
     } specialization;
 
+    TargetConfig target;
+
     struct {
         struct { void* uptr; void (*fn)(void*, String, Module*); } after_pass;
     } hooks;
diff --git a/samples/checkerboard/checkerboard.c b/samples/checkerboard/checkerboard.c
@@ -65,7 +65,7 @@ int main(int argc, char **argv)
 
     info_print("Device-side address is: %zu\n", buf_addr);
 
-    IrArena* a = new_ir_arena(default_arena_config());
+    IrArena* a = new_ir_arena(default_arena_config(&compiler_config.target));
     Module* m;
     if (driver_load_source_file(&compiler_config, SrcSlim, sizeof(checkerboard_kernel_src), checkerboard_kernel_src, "checkerboard", &m) != NoError)
         error("Failed to load checkerboard module");
diff --git a/src/driver/cli.c b/src/driver/cli.c
@@ -95,6 +95,18 @@ F(config->optimisations.inline_everything, inline-everything) \
 F(config->hacks.restructure_everything, restructure-everything) \
 F(config->hacks.recover_structure, recover-structure) \
 
+static IntSizes parse_int_size(String argv) {
+    if (strcmp(argv, "8") == 0)
+        return IntTy8;
+    if (strcmp(argv, "16") == 0)
+        return IntTy16;
+    if (strcmp(argv, "32") == 0)
+        return IntTy32;
+    if (strcmp(argv, "64") == 0)
+        return IntTy64;
+    error("Valid pointer sizes are 8, 16, 32 or 64.");
+}
+
 void cli_parse_compiler_config_args(CompilerConfig* config, int* pargc, char** argv) {
     int argc = *pargc;
 
@@ -142,6 +154,14 @@ void cli_parse_compiler_config_args(CompilerConfig* config, int* pargc, char** a
                 default: break;
             }
             config->specialization.execution_model = em;
+        } else if (strcmp(argv[i], "--word-size") == 0) {
+            argv[i] = NULL;
+            i++;
+            config->target.memory.word_size = parse_int_size(argv[i]);
+        } else if (strcmp(argv[i], "--pointer-size") == 0) {
+            argv[i] = NULL;
+            i++;
+            config->target.memory.ptr_size = parse_int_size(argv[i]);
         } else if (strcmp(argv[i], "--help") == 0 || strcmp(argv[i], "-h") == 0) {
             help = true;
             continue;
@@ -157,6 +177,8 @@ void cli_parse_compiler_config_args(CompilerConfig* config, int* pargc, char** a
         error_print("  --no-dynamic-scheduling                   Disable the built-in dynamic scheduler, restricts code to only leaf functions\n");
         error_print("  --simt2d                                  Emits SIMD code instead of SIMT, only effective with the C backend.\n");
         error_print("  --entry-point <foo>                       Selects an entry point for the program to be specialized on.\n");
+        error_print("  --word-size <8|16|32|64>                  Sets the word size for physical memory emulation (default=32)\n");
+        error_print("  --pointer-size <8|16|32|64>               Sets the pointer size for physical pointers (default=64)\n");
 #define EM(name, _) #name", "
         error_print("  --execution-model <em>                   Selects an entry point for the program to be specialized on.\nPossible values: " EXECUTION_MODELS(EM));
 #undef EM
diff --git a/src/driver/driver.c b/src/driver/driver.c
@@ -50,7 +50,7 @@ ShadyErrorCodes driver_load_source_file(const CompilerConfig* config, SourceLang
         }
         case SrcSPIRV: {
 #ifdef SPV_PARSER_PRESENT
-            parse_spirv_into_shady(len, file_contents, name, mod);
+            parse_spirv_into_shady(config, len, file_contents, name, mod);
 #else
             assert(false && "SPIR-V front-end missing in this version");
 #endif
diff --git a/src/driver/slim.c b/src/driver/slim.c
@@ -22,7 +22,7 @@ int main(int argc, char** argv) {
     cli_parse_compiler_config_args(&args.config, &argc, argv);
     cli_parse_input_files(args.input_filenames, &argc, argv);
 
-    IrArena* arena = new_ir_arena(default_arena_config());
+    IrArena* arena = new_ir_arena(default_arena_config(&args.config.target));
     Module* mod = new_module(arena, "my_module"); // TODO name module after first filename, or perhaps the last one
 
     driver_load_source_files(&args, mod);
diff --git a/src/driver/vcc.c b/src/driver/vcc.c
@@ -68,7 +68,7 @@ int main(int argc, char** argv) {
         exit(MissingInputArg);
     }
 
-    ArenaConfig aconfig = default_arena_config();
+    ArenaConfig aconfig = default_arena_config(&args.config.target);
     IrArena* arena = new_ir_arena(aconfig);
 
     int clang_retval = system("clang --version");
diff --git a/src/frontends/llvm/l2s.c b/src/frontends/llvm/l2s.c
@@ -248,7 +248,7 @@ bool parse_llvm_into_shady(const CompilerConfig* config, size_t len, const char*
     }
     info_print("LLVM IR parsed successfully\n");
 
-    ArenaConfig aconfig = default_arena_config();
+    ArenaConfig aconfig = default_arena_config(&config->target);
     aconfig.check_types = false;
     aconfig.allow_fold = false;
 
diff --git a/src/frontends/slim/parser.c b/src/frontends/slim/parser.c
@@ -1108,7 +1108,7 @@ static void parse_shady_ir(ParserConfig config, const char* contents, Module* mo
 #include "transform/internal_constants.h"
 
 Module* parse_slim_module(const CompilerConfig* config, ParserConfig pconfig, const char* contents, String name) {
-    ArenaConfig aconfig = default_arena_config();
+    ArenaConfig aconfig = default_arena_config(&config->target);
     aconfig.name_bound = false;
     aconfig.check_op_classes = false;
     aconfig.check_types = false;
diff --git a/src/frontends/spirv/s2s.c b/src/frontends/spirv/s2s.c
@@ -1356,8 +1356,8 @@ bool compare_spvid(SpvId* pa, SpvId* pb) {
     return *pa == *pb;
 }
 
-S2SError parse_spirv_into_shady(size_t len, const char* data, String name, Module** dst) {
-    IrArena* a = new_ir_arena(default_arena_config());
+S2SError parse_spirv_into_shady(const CompilerConfig* config, size_t len, const char* data, String name, Module** dst) {
+    IrArena* a = new_ir_arena(default_arena_config(&config->target));
     *dst = new_module(a, name);
 
     SpvParser parser = {
diff --git a/src/frontends/spirv/s2s.h b/src/frontends/spirv/s2s.h
@@ -8,6 +8,6 @@ typedef enum {
     S2S_FailedParsingGeneric,
 } S2SError;
 
-S2SError parse_spirv_into_shady(size_t len, const char* data, String name, Module**);
+S2SError parse_spirv_into_shady(const CompilerConfig* config, size_t len, const char* data, String name, Module**);
 
 #endif
diff --git a/src/runtime/runtime_test.c b/src/runtime/runtime_test.c
@@ -48,7 +48,7 @@ int main(int argc, char* argv[]) {
 
     Program* program;
     IrArena* arena = NULL;
-    arena = new_ir_arena(default_arena_config());
+    arena = new_ir_arena(default_arena_config(&args.driver_config.config.target));
     if (entries_count_list(args.driver_config.input_filenames) == 0) {
         Module* module;
         driver_load_source_file(&args.driver_config.config, SrcSlim, strlen(default_shader), default_shader, "runtime_test", &module);
diff --git a/src/shady/CMakeLists.txt b/src/shady/CMakeLists.txt
@@ -37,6 +37,7 @@ set(SHADY_SOURCES
     compile.c
     annotation.c
     module.c
+    config.c
 
     analysis/cfg.c
     analysis/cfg_dump.c
diff --git a/src/shady/compile.c b/src/shady/compile.c
@@ -10,73 +10,6 @@
 
 #include <stdbool.h>
 
-#define KiB * 1024
-#define MiB * 1024 KiB
-
-CompilerConfig default_compiler_config() {
-    return (CompilerConfig) {
-        .dynamic_scheduling = true,
-        .per_thread_stack_size = 4 KiB,
-
-        .target_spirv_version = {
-            .major = 1,
-            .minor = 4
-        },
-
-        .lower = {
-            .emulate_physical_memory = true,
-            .emulate_generic_ptrs = true,
-        },
-
-        .logging = {
-            // most of the time, we are not interested in seeing generated & internal code in the debug output
-            //.print_internal = true,
-            //.print_generated = true,
-            .print_builtin = true,
-        },
-
-        .optimisations = {
-            .cleanup = {
-                .after_every_pass = true,
-                .delete_unused_instructions = true,
-            }
-        },
-
-        .specialization = {
-            .subgroup_size = 8,
-            .entry_point = NULL
-        }
-    };
-}
-
-ArenaConfig default_arena_config() {
-    ArenaConfig config = {
-        .is_simt = true,
-        .name_bound = true,
-        .allow_fold = true,
-        .check_types = true,
-        .validate_builtin_types = true,
-        .check_op_classes = true,
-
-        .memory = {
-            .word_size = IntTy8,
-            .ptr_size = IntTy64,
-        },
-
-        .optimisations = {
-            .delete_unreachable_structured_cases = true,
-        },
-    };
-
-    for (size_t i = 0; i < NumAddressSpaces; i++) {
-        // by default, all address spaces are physical !
-        config.address_spaces[i].physical = true;
-        config.address_spaces[i].allowed = true;
-    }
-
-    return config;
-}
-
 void add_scheduler_source(const CompilerConfig* config, Module* dst) {
     ParserConfig pconfig = {
         .front_end = true,
diff --git a/src/shady/config.c b/src/shady/config.c
@@ -0,0 +1,76 @@
+#include "shady/ir.h"
+
+#define KiB * 1024
+#define MiB * 1024 KiB
+
+CompilerConfig default_compiler_config() {
+    return (CompilerConfig) {
+        .dynamic_scheduling = true,
+        .per_thread_stack_size = 4 KiB,
+
+        .target_spirv_version = {
+            .major = 1,
+            .minor = 4
+        },
+
+        .lower = {
+            .emulate_physical_memory = true,
+            .emulate_generic_ptrs = true,
+        },
+
+        .logging = {
+            // most of the time, we are not interested in seeing generated & internal code in the debug output
+            //.print_internal = true,
+            //.print_generated = true,
+            .print_builtin = true,
+        },
+
+        .optimisations = {
+            .cleanup = {
+                .after_every_pass = true,
+                .delete_unused_instructions = true,
+            }
+        },
+
+        .target = default_target_config(),
+
+        .specialization = {
+            .subgroup_size = 8,
+            .entry_point = NULL
+        }
+    };
+}
+
+TargetConfig default_target_config() {
+    return (TargetConfig) {
+        .memory = {
+            .word_size = IntTy32,
+            .ptr_size = IntTy64,
+        },
+    };
+}
+
+ArenaConfig default_arena_config(const TargetConfig* target) {
+    ArenaConfig config = {
+        .is_simt = true,
+        .name_bound = true,
+        .allow_fold = true,
+        .check_types = true,
+        .validate_builtin_types = true,
+        .check_op_classes = true,
+
+        .optimisations = {
+            .delete_unreachable_structured_cases = true,
+        },
+
+        .memory = target->memory
+    };
+
+    for (size_t i = 0; i < NumAddressSpaces; i++) {
+        // by default, all address spaces are physical !
+        config.address_spaces[i].physical = true;
+        config.address_spaces[i].allowed = true;
+    }
+
+    return config;
+}
diff --git a/test/test_math.c b/test/test_math.c
@@ -91,7 +91,8 @@ static void test_int_literals(IrArena* a) {
 int main(int argc, char** argv) {
     cli_parse_common_args(&argc, argv);
 
-    ArenaConfig acfg = default_arena_config();
+    TargetConfig target_config = default_target_config();
+    ArenaConfig acfg = default_arena_config(&target_config);
     acfg.check_types = true;
     acfg.allow_fold = true;
     IrArena* a = new_ir_arena(acfg);

Original file line number	Diff line number	Diff line change
`@@ -50,7 +50,7 @@ ShadyErrorCodes driver_load_source_file(const CompilerConfig* config, SourceLang`
`50`	`50`	`}`
`51`	`51`	`case SrcSPIRV: {`
`52`	`52`	`#ifdef SPV_PARSER_PRESENT`
`53`		`- parse_spirv_into_shady(len, file_contents, name, mod);`
	`53`	`+ parse_spirv_into_shady(config, len, file_contents, name, mod);`
`54`	`54`	`#else`
`55`	`55`	`assert(false && "SPIR-V front-end missing in this version");`
`56`	`56`	`#endif`
Original file line number	Diff line number	Diff line change
`@@ -68,7 +68,7 @@ int main(int argc, char** argv) {`
`68`	`68`	`exit(MissingInputArg);`
`69`	`69`	`}`
`70`	`70`
`71`		`- ArenaConfig aconfig = default_arena_config();`
	`71`	`+ ArenaConfig aconfig = default_arena_config(&args.config.target);`
`72`	`72`	`IrArena* arena = new_ir_arena(aconfig);`
`73`	`73`
`74`	`74`	`int clang_retval = system("clang --version");`
Original file line number	Diff line number	Diff line change
`@@ -248,7 +248,7 @@ bool parse_llvm_into_shady(const CompilerConfig* config, size_t len, const char*`
`248`	`248`	`}`
`249`	`249`	`info_print("LLVM IR parsed successfully\n");`
`250`	`250`
`251`		`- ArenaConfig aconfig = default_arena_config();`
	`251`	`+ ArenaConfig aconfig = default_arena_config(&config->target);`
`252`	`252`	`aconfig.check_types = false;`
`253`	`253`	`aconfig.allow_fold = false;`
`254`	`254`
Original file line number	Diff line number	Diff line change
`@@ -1356,8 +1356,8 @@ bool compare_spvid(SpvId* pa, SpvId* pb) {`
`1356`	`1356`	`return pa == pb;`
`1357`	`1357`	`}`
`1358`	`1358`
`1359`		`-S2SError parse_spirv_into_shady(size_t len, const char* data, String name, Module** dst) {`
`1360`		`- IrArena* a = new_ir_arena(default_arena_config());`
	`1359`	`+S2SError parse_spirv_into_shady(const CompilerConfig* config, size_t len, const char* data, String name, Module** dst) {`
	`1360`	`+ IrArena* a = new_ir_arena(default_arena_config(&config->target));`
`1361`	`1361`	`*dst = new_module(a, name);`
`1362`	`1362`
`1363`	`1363`	`SpvParser parser = {`