Make it compile on ROCm

Author: Dmitry Kozlov

RadeonRays/src/kernels/CL/intersect_bvh2_lds.cl

@@ -75,6 +75,8 @@ KERNEL void intersect_main(
     // Hit data
     GLOBAL Intersection *hits)
 {
+    __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
+
     uint index = get_global_id(0);
     uint local_index = get_local_id(0);
 
@@ -85,8 +87,6 @@ KERNEL void intersect_main(
 
         if (ray_is_active(&my_ray))
         {
-            __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
-
             const float3 invDir = safe_invdir(my_ray);
             const float3 oxInvDir = -my_ray.o.xyz * invDir;
 
@@ -231,6 +231,8 @@ KERNEL void occluded_main(
     // Hit results: 1 for hit and -1 for miss
     GLOBAL int *hits)
 {
+    __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
+
     uint index = get_global_id(0);
     uint local_index = get_local_id(0);
 
@@ -241,8 +243,6 @@ KERNEL void occluded_main(
 
         if (ray_is_active(&my_ray))
         {
-            __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
-
             const float3 invDir = safe_invdir(my_ray);
             const float3 oxInvDir = -my_ray.o.xyz * invDir;
 

RadeonRays/src/kernels/CL/intersect_bvh2_lds_fp16.cl

@@ -140,6 +140,8 @@ KERNEL void intersect_main(
     // Hit data
     GLOBAL Intersection *hits)
 {
+    __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
+
     uint index = get_global_id(0);
     uint local_index = get_local_id(0);
 
@@ -150,8 +152,6 @@ KERNEL void intersect_main(
 
         if (ray_is_active(&my_ray))
         {
-            __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
-
             // Precompute inverse direction and origin / dir for bbox testing
             const float3 invDir32 = safe_invdir2(my_ray);
             const half3 invDir = convert_half3(invDir32);
@@ -319,6 +319,8 @@ KERNEL void occluded_main(
     // Hit results: 1 for hit and -1 for miss
     GLOBAL int *hits)
 {
+    __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
+
     uint index = get_global_id(0);
     uint local_index = get_local_id(0);
 
@@ -329,8 +331,6 @@ KERNEL void occluded_main(
 
         if (ray_is_active(&my_ray))
         {
-            __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
-
             // Precompute inverse direction and origin / dir for bbox testing
             const float3 invDir32 = safe_invdir2(my_ray);
             const half3 invDir = convert_half3(invDir32);

RadeonRays/src/kernels/CL/intersect_bvh2_short_stack.cl

@@ -142,6 +142,9 @@ occluded_main(
     GLOBAL int* hits
     )
 {
+    // Allocate stack in LDS
+    __local int lds[SHORT_STACK_SIZE * WAVEFRONT_SIZE];
+
     int global_id = get_global_id(0);
     int local_id = get_local_id(0);
     int group_id = get_group_id(0);
@@ -156,8 +159,7 @@ occluded_main(
             // Allocate stack in global memory 
             __global int* gm_stack_base = stack + (group_id * WAVEFRONT_SIZE + local_id) * GLOBAL_STACK_SIZE;
             __global int* gm_stack = gm_stack_base;
-            // Allocate stack in LDS
-            __local int lds[SHORT_STACK_SIZE * WAVEFRONT_SIZE];
+
             __local int* lm_stack_base = lds + local_id;
             __local int* lm_stack = lm_stack_base;
 
@@ -293,6 +295,9 @@ KERNEL void intersect_main(
     // Hit data
     GLOBAL Intersection* hits)
 {
+    // Allocate stack in LDS
+    __local int lds[SHORT_STACK_SIZE * WAVEFRONT_SIZE];
+
     int global_id = get_global_id(0);
     int local_id = get_local_id(0);
     int group_id = get_group_id(0);
@@ -307,8 +312,6 @@ KERNEL void intersect_main(
             // Allocate stack in global memory 
             __global int* gm_stack_base = stack + (group_id * WAVEFRONT_SIZE + local_id) * GLOBAL_STACK_SIZE;
             __global int* gm_stack = gm_stack_base;
-            // Allocate stack in LDS
-            __local int lds[SHORT_STACK_SIZE * WAVEFRONT_SIZE];
             __local int* lm_stack_base = lds + local_id;
             __local int* lm_stack = lm_stack_base;