Skip to content

Commit bf0a1dc

Browse files
Hugobros3Hugobros3
authored
Merge branch 'master' into fix-double-free
2 parents 75f4a3e + 20827a2 commit bf0a1dc

File tree

105 files changed

+1630
-1286
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

105 files changed

+1630
-1286
lines changed

‎CMakeLists.txt

Lines changed: 7 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -8,11 +8,18 @@ find_package(Git)
88

99
set(CMAKE_C_STANDARD 11)
1010
set(CMAKE_C_STANDARD_REQUIRED ON)
11+
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
12+
13+
option(BUILD_SHARED_LIBS "Build using shared libraries" ON)
1114

1215
if (MSVC)
1316
add_compile_definitions(_CRT_SECURE_NO_WARNINGS)
1417
endif()
1518

19+
if (WIN32)
20+
set(BUILD_SHARED_LIBS OFF)
21+
endif()
22+
1623
add_subdirectory(SPIRV-Headers)
1724

1825
if(CMAKE_VERSION VERSION_GREATER_EQUAL "3.24")
@@ -98,20 +105,11 @@ endif()
98105

99106
include(CMakePackageConfigHelpers)
100107

101-
install(TARGETS api EXPORT shady_export_set)
102-
install(TARGETS shady EXPORT shady_export_set ARCHIVE DESTINATION ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
103-
104108
if (TARGET vcc)
105109
add_subdirectory(vcc-std)
106110
endif ()
107111

108-
if (TARGET runtime)
109-
install(TARGETS runtime EXPORT shady_export_set ARCHIVE DESTINATION ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
110-
endif()
111-
112-
install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/shady DESTINATION include)
113112
install(EXPORT shady_export_set DESTINATION share/cmake/shady/ NAMESPACE shady:: FILE shady-targets.cmake)
114113

115114
configure_file(cmake/shady-config.cmake.in shady-config.cmake @ONLY)
116115
install(FILES "${CMAKE_CURRENT_BINARY_DIR}/shady-config.cmake" DESTINATION share/cmake/shady)
117-
#install(FILES "${CMAKE_CURRENT_BINARY_DIR}/shady-config.cmake" DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/shady)

‎include/shady/driver.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,8 +26,8 @@ typedef enum {
2626
} SourceLanguage;
2727

2828
SourceLanguage guess_source_language(const char* filename);
29-
ShadyErrorCodes driver_load_source_file(SourceLanguage lang, size_t, const char* file_contents, Module* mod);
30-
ShadyErrorCodes driver_load_source_file_from_filename(const char* filename, Module* mod);
29+
ShadyErrorCodes driver_load_source_file(const CompilerConfig* config, SourceLanguage lang, size_t, const char* file_contents, String, Module** mod);
30+
ShadyErrorCodes driver_load_source_file_from_filename(const CompilerConfig* config, const char* filename, String, Module** mod);
3131

3232
typedef enum {
3333
TgtAuto,

‎include/shady/grammar.json

Lines changed: 13 additions & 33 deletions
Original file line numberDiff line numberDiff line change
@@ -2,53 +2,32 @@
22
"address-spaces": [
33
{
44
"name": "Generic",
5-
"llvm-id": 0,
6-
"physical": true
5+
"llvm-id": 0
76
},
87
{
9-
"name": "GlobalPhysical",
8+
"name": "Global",
109
"description": "Global memory, all threads see the same data (not necessarily consistent!)",
11-
"llvm-id": 1,
12-
"physical": true
10+
"llvm-id": 1
1311
},
1412
{
15-
"name": "SharedPhysical",
13+
"name": "Shared",
1614
"description": "Points into workgroup-private memory (aka shared memory)",
17-
"llvm-id": 3,
18-
"physical": true
15+
"llvm-id": 3
1916
},
2017
{
21-
"name": "SubgroupPhysical",
18+
"name": "Subgroup",
2219
"description": [
2320
"Points into subgroup-private memory",
2421
"All threads in a subgroup see the same contents for the same address, but threads in different subgroups see different data.",
2522
"Needs to be lowered to something else since targets do not understand this" ],
26-
"llvm-id": 9,
27-
"physical": true
23+
"llvm-id": 9
2824
},
2925
{
30-
"name": "PrivatePhysical",
26+
"name": "Private",
3127
"description": [
3228
"Points into thread-private memory (all threads see different contents for the same address)"
3329
],
34-
"llvm-id": 5,
35-
"physical": true
36-
},
37-
{
38-
"name": "GlobalLogical",
39-
"llvm-id": 388
40-
},
41-
{
42-
"name": "SharedLogical",
43-
"llvm-id": 387
44-
},
45-
{
46-
"name": "SubgroupLogical",
47-
"llvm-id": 386
48-
},
49-
{
50-
"name": "PrivateLogical",
51-
"llvm-id": 385
30+
"llvm-id": 5
5231
},
5332
{
5433
"name": "Input",
@@ -74,8 +53,8 @@
7453
"llvm-id": 392
7554
},
7655
{
77-
"name": "FunctionLogical",
78-
"description": "Weird SPIR-V nonsense: this is like PrivateLogical, but with non-static lifetimes (ie function lifetime)",
56+
"name": "Function",
57+
"description": "Weird SPIR-V nonsense: this is like Private, but with non-static lifetimes (ie function lifetime)",
7958
"llvm-id": 393
8059
},
8160
{
@@ -225,7 +204,8 @@
225204
"class": "type",
226205
"ops": [
227206
{ "name": "address_space", "type": "AddressSpace" },
228-
{ "name": "pointed_type", "class": "type" }
207+
{ "name": "pointed_type", "class": "type" },
208+
{ "name": "is_reference", "type": "bool" }
229209
]
230210
},
231211
{

‎include/shady/ir.h

Lines changed: 9 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -69,12 +69,13 @@ typedef struct {
6969
bool check_op_classes;
7070
bool check_types;
7171
bool allow_fold;
72-
bool untyped_ptrs;
7372
bool validate_builtin_types; // do @Builtins variables need to match their type in builtins.h ?
7473
bool is_simt;
7574

76-
bool allow_subgroup_memory;
77-
bool allow_shared_memory;
75+
struct {
76+
bool physical;
77+
bool allowed;
78+
} address_spaces[NumAddressSpaces];
7879

7980
struct {
8081
/// Selects which type the subgroup intrinsic primops use to manipulate masks
@@ -101,7 +102,6 @@ typedef struct {
101102
} optimisations;
102103
} ArenaConfig;
103104

104-
typedef struct CompilerConfig_ CompilerConfig;
105105
ArenaConfig default_arena_config();
106106

107107
IrArena* new_ir_arena(ArenaConfig);
@@ -139,6 +139,7 @@ Nodes filter_out_annotation(IrArena*, Nodes, const char* name);
139139

140140
bool is_abstraction (const Node*);
141141
String get_abstraction_name (const Node* abs);
142+
String get_abstraction_name_unsafe(const Node* abs);
142143
const Node* get_abstraction_body (const Node* abs);
143144
Nodes get_abstraction_params(const Node* abs);
144145

@@ -255,7 +256,7 @@ const Node* bind_last_instruction_and_wrap_in_block(BodyBuilder*, const Node*);
255256

256257
//////////////////////////////// Compilation ////////////////////////////////
257258

258-
struct CompilerConfig_ {
259+
typedef struct CompilerConfig_ {
259260
bool dynamic_scheduling;
260261
uint32_t per_thread_stack_size;
261262

@@ -278,8 +279,8 @@ struct CompilerConfig_ {
278279
struct {
279280
bool spv_shuffle_instead_of_broadcast_first;
280281
bool force_join_point_lifting;
281-
bool assume_no_physical_global_ptrs;
282282
bool restructure_everything;
283+
bool recover_structure;
283284
} hacks;
284285

285286
struct {
@@ -315,7 +316,7 @@ struct CompilerConfig_ {
315316
struct {
316317
struct { void* uptr; void (*fn)(void*, String, Module*); } after_pass;
317318
} hooks;
318-
};
319+
} CompilerConfig;
319320

320321
CompilerConfig default_compiler_config();
321322

@@ -324,6 +325,7 @@ typedef enum CompilationResult_ {
324325
} CompilationResult;
325326

326327
CompilationResult run_compiler_passes(CompilerConfig* config, Module** mod);
328+
void link_module(Module* dst, Module* src);
327329

328330
//////////////////////////////// Emission ////////////////////////////////
329331

‎include/shady/runtime.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,6 @@ typedef struct CompilerConfig_ CompilerConfig;
2929
typedef struct Module_ Module;
3030

3131
Program* new_program_from_module(Runtime*, const CompilerConfig*, Module*);
32-
Program* load_program(Runtime*, const CompilerConfig*, const char* program_src);
33-
Program* load_program_from_disk(Runtime*, const CompilerConfig*, const char* path);
3432

3533
Command* launch_kernel(Program*, Device*, const char* entry_point, int dimx, int dimy, int dimz, int args_count, void** args);
3634
bool wait_completion(Command*);

‎samples/aobench/CMakeLists.txt

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
if (TARGET vcc)
22
add_executable(aobench_host ao_host.c ao_main.c)
3-
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)
3+
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 -O3 -o ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll COMMENT ao.comp.c.ll DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/ao.comp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/ao.c)
44

55
set_property(SOURCE ${CMAKE_CURRENT_SOURCE_DIR}/ao_main.c APPEND PROPERTY OBJECT_DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll)
66
add_custom_command(TARGET aobench_host POST_BUILD
@@ -11,7 +11,7 @@ if (TARGET vcc)
1111
find_program(ISPC_EXE "ispc")
1212
if (ISPC_EXE)
1313
target_compile_definitions(aobench_host PUBLIC ENABLE_ISPC=1)
14-
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)
14+
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 --restructure-everything COMMENT generating aobench.ispc DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/ao.comp.c.ll)
1515
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)
1616
add_library(aobench_ispc OBJECT ${CMAKE_CURRENT_BINARY_DIR}/aobench.ispc.o)
1717
set_target_properties(aobench_ispc PROPERTIES LINKER_LANGUAGE C)

‎samples/aobench/ao.c

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
#define FUNCTION
77
#endif
88

9-
FUNCTION static unsigned int FNVHash(char* str, unsigned int length) {
9+
FUNCTION unsigned int FNVHash(char* str, unsigned int length) {
1010
const unsigned int fnv_prime = 0x811C9DC5;
1111
unsigned int hash = 0;
1212
unsigned int i = 0;
@@ -20,31 +20,31 @@ FUNCTION static unsigned int FNVHash(char* str, unsigned int length) {
2020
return hash;
2121
}
2222

23-
FUNCTION static unsigned int nrand(unsigned int* rng) {
23+
FUNCTION unsigned int nrand(unsigned int* rng) {
2424
unsigned int orand = *rng;
2525
*rng = FNVHash((char*) &orand, 4);
2626
return *rng;
2727
}
2828

29-
FUNCTION static Scalar drand48(Ctx* ctx) {
29+
FUNCTION Scalar drand48(Ctx* ctx) {
3030
Scalar n = (nrand(&ctx->rng) / 65536.0f);
3131
n = n - floorf(n);
3232
return n;
3333
}
3434

35-
FUNCTION static Scalar vdot(vec v0, vec v1)
35+
FUNCTION Scalar vdot(vec v0, vec v1)
3636
{
3737
return v0.x * v1.x + v0.y * v1.y + v0.z * v1.z;
3838
}
3939

40-
FUNCTION static void vcross(vec *c, vec v0, vec v1)
40+
FUNCTION void vcross(vec *c, vec v0, vec v1)
4141
{
4242
c->x = v0.y * v1.z - v0.z * v1.y;
4343
c->y = v0.z * v1.x - v0.x * v1.z;
4444
c->z = v0.x * v1.y - v0.y * v1.x;
4545
}
4646

47-
FUNCTION static void vnormalize(vec *c)
47+
FUNCTION void vnormalize(vec *c)
4848
{
4949
Scalar length = sqrtf(vdot((*c), (*c)));
5050

@@ -55,7 +55,7 @@ FUNCTION static void vnormalize(vec *c)
5555
}
5656
}
5757

58-
FUNCTION static void
58+
FUNCTION void
5959
ray_sphere_intersect(Isect *isect, const Ray *ray, const Sphere *sphere)
6060
{
6161
vec rs = { 0 };
@@ -88,7 +88,7 @@ ray_sphere_intersect(Isect *isect, const Ray *ray, const Sphere *sphere)
8888
}
8989
}
9090

91-
FUNCTION static void
91+
FUNCTION void
9292
ray_plane_intersect(Isect *isect, const Ray *ray, const Plane *plane)
9393
{
9494
Scalar d = -vdot(plane->p, plane->n);
@@ -110,7 +110,7 @@ ray_plane_intersect(Isect *isect, const Ray *ray, const Plane *plane)
110110
}
111111
}
112112

113-
FUNCTION static void
113+
FUNCTION void
114114
orthoBasis(vec *basis, vec n)
115115
{
116116
basis[2] = n;
@@ -133,7 +133,7 @@ orthoBasis(vec *basis, vec n)
133133
vnormalize(&basis[1]);
134134
}
135135

136-
FUNCTION static void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
136+
FUNCTION void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
137137
{
138138
int i, j;
139139
int ntheta = NAO_SAMPLES;
@@ -193,7 +193,7 @@ FUNCTION static void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
193193
col->z = occlusion;
194194
}
195195

196-
FUNCTION static unsigned char aobench_clamp(Scalar f)
196+
FUNCTION unsigned char aobench_clamp(Scalar f)
197197
{
198198
Scalar s = (f * 255.5f);
199199

@@ -203,13 +203,13 @@ FUNCTION static unsigned char aobench_clamp(Scalar f)
203203
return (unsigned char) s;
204204
}
205205

206-
FUNCTION EXTERNAL_FN Ctx get_init_context() {
206+
EXTERNAL_FN Ctx get_init_context() {
207207
return (Ctx) {
208208
.rng = 0xFEEFDEED,
209209
};
210210
}
211211

212-
FUNCTION EXTERNAL_FN void render_pixel(Ctx* ctx, int x, int y, int w, int h, int nsubsamples, unsigned char* img) {
212+
EXTERNAL_FN void render_pixel(Ctx* ctx, int x, int y, int w, int h, int nsubsamples, TEXEL_T* img) {
213213
Scalar pixel[3] = { 0, 0, 0 };
214214

215215
ctx->rng = x * w + y;
@@ -269,7 +269,7 @@ FUNCTION EXTERNAL_FN void render_pixel(Ctx* ctx, int x, int y, int w, int h, int
269269
img[3 * (y * w + x) + 2] = aobench_clamp(pixel[2]);
270270
}
271271

272-
FUNCTION EXTERNAL_FN void init_scene(Ctx* ctx)
272+
EXTERNAL_FN void init_scene(Ctx* ctx)
273273
{
274274
ctx->spheres[0].center.x = -2.0f;
275275
ctx->spheres[0].center.y = 0.0f;
@@ -294,4 +294,4 @@ FUNCTION EXTERNAL_FN void init_scene(Ctx* ctx)
294294
ctx->plane.n.y = 1.0f;
295295
ctx->plane.n.z = 0.0f;
296296

297-
}
297+
}

‎samples/aobench/ao.comp.cpp

Lines changed: 10 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -32,27 +32,21 @@ float cosf(float) __asm__("shady::prim_op::cos");
3232
float fmodf(float, float) __asm__("shady::prim_op::mod");
3333
float fabsf(float) __asm__("shady::prim_op::abs");
3434
float floorf(float) __asm__("shady::prim_op::floor");
35+
36+
#define EXTERNAL_FN static
37+
#define FUNCTION static
38+
3539
#include "ao.c"
3640

37-
extern "C" __attribute__((annotate("shady::workgroup_size::16::16::1")))
38-
compute_shader void aobench_kernel(global unsigned char* out) {
39-
//outColor = (vec4) { fragColor[0], fragColor[1], fragColor[2], 1.0f };
40-
//outColor = (vec4) { fragCoord[0] / 1024, fragCoord[1] / 1024, 1.0f, 1.0f };
41+
#define xstr(s) str(s)
42+
#define str(s) #s
4143

44+
extern "C" __attribute__((annotate("shady::workgroup_size::" xstr(BLOCK_SIZE) "::" xstr(BLOCK_SIZE) "::1")))
45+
compute_shader void aobench_kernel(global TEXEL_T* out) {
4246
Ctx ctx = get_init_context();
4347
init_scene(&ctx);
4448

4549
int x = global_id.x;
4650
int y = global_id.y;
47-
//int x = (int) fragCoord.x % 1024;
48-
//int y = (int) fragCoord.y % 1024;
49-
50-
// unsigned int out[3]; // = { 55, 0, 0};
51-
out[0] = 255;
52-
out[1] = 255;
53-
render_pixel(&ctx, x + 3, y, WIDTH, HEIGHT, NSUBSAMPLES, (unsigned char*) out);
54-
//out[2] = 155;
55-
// out[0] = x / 4;
56-
// out[1] = y / 4;
57-
//outColor = (vec4) { ((int) out[0]) / 255.0f, ((int) out[1]) / 255.0f, ((int) out[2]) / 255.0f, 1.0f };
58-
}
51+
render_pixel(&ctx, x, y, WIDTH, HEIGHT, NSUBSAMPLES, (TEXEL_T*) out);
52+
}

0 commit comments

Comments
 (0)