fix ao.cu

Author: Hugobros3

samples/aobench/ao.c

@@ -6,7 +6,7 @@
 #define FUNCTION
 #endif
 
-FUNCTION static unsigned int FNVHash(char* str, unsigned int length) {
+FUNCTION unsigned int FNVHash(char* str, unsigned int length) {
     const unsigned int fnv_prime = 0x811C9DC5;
     unsigned int hash = 0;
     unsigned int i = 0;
@@ -20,31 +20,31 @@ FUNCTION static unsigned int FNVHash(char* str, unsigned int length) {
     return hash;
 }
 
-FUNCTION static unsigned int nrand(unsigned int* rng) {
+FUNCTION unsigned int nrand(unsigned int* rng) {
     unsigned int orand = *rng;
     *rng = FNVHash((char*) &orand, 4);
     return *rng;
 }
 
-FUNCTION static Scalar drand48(Ctx* ctx) {
+FUNCTION Scalar drand48(Ctx* ctx) {
     Scalar n = (nrand(&ctx->rng) / 65536.0f);
     n = n - floorf(n);
     return n;
 }
 
-FUNCTION static Scalar vdot(vec v0, vec v1)
+FUNCTION Scalar vdot(vec v0, vec v1)
 {
     return v0.x * v1.x + v0.y * v1.y + v0.z * v1.z;
 }
 
-FUNCTION static void vcross(vec *c, vec v0, vec v1)
+FUNCTION void vcross(vec *c, vec v0, vec v1)
 {
     c->x = v0.y * v1.z - v0.z * v1.y;
     c->y = v0.z * v1.x - v0.x * v1.z;
     c->z = v0.x * v1.y - v0.y * v1.x;
 }
 
-FUNCTION static void vnormalize(vec *c)
+FUNCTION void vnormalize(vec *c)
 {
     Scalar length = sqrtf(vdot((*c), (*c)));
 
@@ -55,7 +55,7 @@ FUNCTION static void vnormalize(vec *c)
     }
 }
 
-FUNCTION static void
+FUNCTION void
 ray_sphere_intersect(Isect *isect, const Ray *ray, const Sphere *sphere)
 {
     vec rs = { 0 };
@@ -88,7 +88,7 @@ ray_sphere_intersect(Isect *isect, const Ray *ray, const Sphere *sphere)
     }
 }
 
-FUNCTION static void
+FUNCTION void
 ray_plane_intersect(Isect *isect, const Ray *ray, const Plane *plane)
 {
     Scalar d = -vdot(plane->p, plane->n);
@@ -110,7 +110,7 @@ ray_plane_intersect(Isect *isect, const Ray *ray, const Plane *plane)
     }
 }
 
-FUNCTION static void
+FUNCTION void
 orthoBasis(vec *basis, vec n)
 {
     basis[2] = n;
@@ -133,7 +133,7 @@ orthoBasis(vec *basis, vec n)
     vnormalize(&basis[1]);
 }
 
-FUNCTION static void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
+FUNCTION void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
 {
     int    i, j;
     int    ntheta = NAO_SAMPLES;
@@ -193,7 +193,7 @@ FUNCTION static void ambient_occlusion(Ctx* ctx, vec *col, const Isect *isect)
     col->z = occlusion;
 }
 
-FUNCTION static unsigned char aobench_clamp(Scalar f)
+FUNCTION unsigned char aobench_clamp(Scalar f)
 {
     Scalar s = (f * 255.5f);
 

samples/aobench/ao.cu

@@ -1,11 +1,14 @@
-#define EXTERNAL_FN static inline __device__ __attribute__((always_inline))
-#define FUNCTION static inline __device__ __attribute__((always_inline))
+// #define EXTERNAL_FN static inline __device__ __attribute__((always_inline))
+// #define FUNCTION static inline __device__ __attribute__((always_inline))
+
+#define EXTERNAL_FN static __device__
+#define FUNCTION static __device__
 
 #include "ao.c"
 
 extern "C" {
 
-__global__ void aobench_kernel(unsigned TEXEL_T* out) {
+__global__ void aobench_kernel(TEXEL_T* out) {
     int x = threadIdx.x + blockDim.x * blockIdx.x;
     int y = threadIdx.y + blockDim.y * blockIdx.y;
 

src/shady/emit/spirv/emit_spv_instructions.c

@@ -385,7 +385,10 @@ static void emit_primop(Emitter* emitter, FnBuilder fn_builder, BBBuilder bb_bui
                 assert(results_count == 1);
                 results[0] = result;
             } else {
-                error("TODO: OpPtrAccessChain")
+                const Type* target_type = instr->type;
+                SpvId result = spvb_ptr_access_chain(bb_builder, emit_type(emitter, target_type), base, emit_value(emitter, bb_builder, args.nodes[1]), args.count - 2, indices);
+                assert(results_count == 1);
+                results[0] = result;
             }
             return;
         }