cuda: emit is_first/broadcast_first

Author: Hugobros3

src/shady/emit/c/cuda_builtins.cu

@@ -12,3 +12,20 @@ __device__ void __shady_prepare_builtins() {
     GlobalInvocationId.arr[1] = threadIdx.y + blockDim.y * blockIdx.y;
     GlobalInvocationId.arr[2] = threadIdx.z + blockDim.z * blockIdx.z;
 }
+
+__device__ bool __shady_elect_first() {
+    unsigned int writemask = __activemask();
+    // Find the lowest-numbered active lane
+    int elected_lane = __ffs(writemask) - 1;
+    return threadIdx.x == __shfl_sync(writemask, threadIdx.x, elected_lane)
+        && threadIdx.y == __shfl_sync(writemask, threadIdx.y, elected_lane)
+        && threadIdx.z == __shfl_sync(writemask, threadIdx.z, elected_lane);
+}
+
+template<typename T>
+__device__ T __shady_broadcast_first(T t) {
+    unsigned int writemask = __activemask();
+    // Find the lowest-numbered active lane
+    int elected_lane = __ffs(writemask) - 1;
+    return __shfl_sync(writemask, t, elected_lane);
+}

src/shady/emit/c/emit_c_instructions.c

@@ -618,13 +618,13 @@ static void emit_primop(Emitter* emitter, Printer* p, const Node* node, Instruct
         case get_stack_base_op:
         case push_stack_op:
         case pop_stack_op:
-        case get_stack_pointer_op:
-        case set_stack_pointer_op: error("Stack operations need to be lowered.");
+        case get_stack_size_op:
+        case set_stack_size_op: error("Stack operations need to be lowered.");
         case default_join_point_op:
         case create_joint_point_op: error("lowered in lower_tailcalls.c");
         case subgroup_elect_first_op: {
             switch (emitter->config.dialect) {
-                case CDialect_CUDA: error("TODO")
+                case CDialect_CUDA: term = term_from_cvalue(format_string_arena(emitter->arena->arena, "__shady_elect_first()")); break;
                 case CDialect_ISPC: term = term_from_cvalue(format_string_arena(emitter->arena->arena, "(programIndex == count_trailing_zeros(lanemask()))")); break;
                 case CDialect_C11:
                 case CDialect_GLSL: error("TODO")
@@ -641,7 +641,7 @@ static void emit_primop(Emitter* emitter, Printer* p, const Node* node, Instruct
         case subgroup_broadcast_first_op: {
             CValue value = to_cvalue(emitter, emit_value(emitter, p, first(prim_op->operands)));
             switch (emitter->config.dialect) {
-                case CDialect_CUDA: error("TODO")
+                case CDialect_CUDA: term = term_from_cvalue(format_string_arena(emitter->arena->arena, "__shady_broadcast_first(%s)", value)); break;
                 case CDialect_ISPC: term = term_from_cvalue(format_string_arena(emitter->arena->arena, "extract(%s, count_trailing_zeros(lanemask()))", value)); break;
                 case CDialect_C11:
                 case CDialect_GLSL: error("TODO")