Logic_kernelv1.cl

/// Enter your kernel in this window__kernel void logicsimulation(uint length,
__kernel void logicsimulationRobust(
    ulong length, __global uint *top_order, __global uint *reverse_top_order,
    __global uint *gate_type, __global uint2 *edges, __global uint2 *out_edges,
    __global uint *p1, __global uint *p2, __global uint *robust_result1,
    __global uint *robust_result2, __global uint *is_robust_path_to_po) {
  unsigned int output_offset = length * get_global_id(0);
  ulong size = length;
  unsigned int result1 = 0;
  unsigned int result2 = 0;
  // int pattern_counter = 0;
  unsigned int pattern1_1;
  unsigned int pattern1_2;
  unsigned int pattern2_1;
  unsigned int pattern2_2;
  unsigned int isRobust_1;
  unsigned int isRobust_2;
  for (int i = 0; i < size; i++) {
    unsigned int c_node = top_order[i];
    result1 = 0;
    result2 = 0;
    unsigned int g = gate_type[c_node];
    if (g) // All other gates, if node is an input the else is executed, the
           // gate type can also be used to determine the in_degree
    {

      pattern1_1 = p1[output_offset + edges[c_node].x];
      pattern1_2 = p1[output_offset + edges[c_node].y];
      pattern2_1 = p2[output_offset + edges[c_node].x];
      pattern2_2 = p2[output_offset + edges[c_node].y];

      if (g == 1) {

        result1 = ~(pattern1_1 & pattern1_2);
        result2 = ~(pattern2_1 & pattern2_2);
      } else if (g == 2) {

        result1 = ~(pattern1_1 | pattern1_2);
        result2 = ~(pattern2_1 | pattern2_2);
      } else if (g == 4) {

        result1 = pattern1_1 & pattern1_2;
        result2 = pattern2_1 & pattern2_2;
      } else if (g == 5) {

        result1 = pattern1_1 | pattern1_2;
        result2 = pattern2_1 | pattern2_2;
      } else if (g == 6) {
        result1 = ~pattern1_1;
        result2 = ~pattern2_1;
      }

      if (g == 1 || g == 4) // AND NAND case
      {
        // isRobustAND NAND
        isRobust_1 = ((~pattern1_1) & pattern1_2 & pattern2_2) |
                     (pattern1_1 & (~pattern1_2) & pattern2_1 & pattern2_2);
        isRobust_2 = ((~pattern2_1) & pattern2_2 & pattern1_2) |
                     (pattern2_1 & (~pattern2_2) & pattern1_1 & pattern1_2);

        //	isRobustOrNonRobust_l = (pattern1_1 ^pattern1_2) & pattern2_2;
        //	isRobustOrNonRobust_2 = (pattern2_1 ^pattern2_2) & pattern1_2;
      } else if (g == 5 || g == 2) // Or Nor
      {
        isRobust_1 =
            (pattern1_1 & (~pattern1_2) & (~pattern2_2)) |
            ((~pattern1_1) & pattern1_2 & (~pattern2_1) & (~pattern2_2));
        isRobust_2 =
            (pattern2_1 & (~pattern2_2) & (~pattern1_2)) |
            ((~pattern2_1) & pattern2_2 & (~pattern1_1) & (~pattern1_2));

        // isRobustOrNonRobust_l = (pattern1_1 ^ pattern1_2) & (~pattern2_2);
        // isRobustOrNonRobust_2 = (pattern2_1 ^ pattern2_2) & (~pattern1_2);
      } else if (g == 6) {
        isRobust_1 = pattern1_1 ^ pattern2_1;
        isRobust_2 = pattern1_1 ^ pattern2_1;

        // isRobustOrNonRobust_l = pattern1_1 ^ pattern2_1;
        //  isRobustOrNonRobust_2 = pattern1_1 ^ pattern2_1;
      }

      p1[output_offset + c_node] = result1;
      p2[output_offset + c_node] = result2;

      robust_result1[output_offset + c_node] = isRobust_1;
      robust_result2[output_offset + c_node] = isRobust_2;
      //	is_robust_or_non_robust_result1[output_offset + c_node] =
      //isRobustOrNonRobust_l; 	is_robust_or_non_robust_result2[output_offset +
      //c_node] = isRobustOrNonRobust_2;

    } else {
      // p1[output_offset + c_node] =
      // boost::dynamic_bitset<>(generate4bit()).to_ulong(); p2[output_offset +
      // c_node] = boost::dynamic_bitset<>(generate4bit()).to_ulong();
    }
  }

  /* traverse over the graph in reverse topological order */
  for (int i = 0; i < size; i++) {
    unsigned int c_node = reverse_top_order[i];
    if ((out_edges[c_node].x != 0) &&
        (out_edges[c_node].y != 0)) // Node is not an output
    {
      unsigned int result = 0;

      unsigned int r1 = 0;
      unsigned int path_po = 0;
      if (edges[out_edges[c_node].x].x == c_node) {
        r1 = robust_result1[output_offset + out_edges[c_node].x];
        path_po = is_robust_path_to_po[output_offset + out_edges[c_node].x];
      } else {
        r1 = robust_result2[output_offset + out_edges[c_node].x];
        path_po = is_robust_path_to_po[output_offset + out_edges[c_node].x];
      }

      result |= (r1 & path_po);

      if (edges[out_edges[c_node].y].x == c_node) {
        r1 = robust_result1[output_offset + out_edges[c_node].y];
        path_po = is_robust_path_to_po[output_offset + out_edges[c_node].y];
      } else {
        r1 = robust_result2[output_offset + out_edges[c_node].y];
        path_po = is_robust_path_to_po[output_offset + out_edges[c_node].y];
      }
      result |= (r1 & path_po);
      is_robust_path_to_po[output_offset + c_node] = result;
    } else {
      is_robust_path_to_po[output_offset + c_node] =
          ~0; // set the bits of all outputs to 1
    }         // Node is an output because no out edges
              // case of output set all bits to "1"
  }

  /* Last Interation over the circuit in topological order*/
  for (int i = 0; i < size; i++) {
    unsigned int c_node = top_order[i];
    unsigned int g = gate_type[c_node];
    if (g) {
      /* Check all input edges */
      unsigned int temp_final_result = 0;
      unsigned int btrofr = 0;
      /* robust_result[output_offset + c_node].x; */

      if (out_edges[edges[c_node].x].x == c_node)
        btrofr = p1[output_offset + edges[c_node].x];
      else
        btrofr = p2[output_offset + edges[c_node].x];

      temp_final_result |=
          is_robust_path_to_po[output_offset + edges[c_node].x] &
          robust_result1[output_offset + c_node] & btrofr;

      /* robust_result[output_offset + c_node].y; */
      if (out_edges[edges[c_node].y].x == c_node)
        btrofr = p1[output_offset + edges[c_node].y];
      else
        btrofr = p2[output_offset + edges[c_node].y];

      temp_final_result |=
          is_robust_path_to_po[output_offset + edges[c_node].y] &
          robust_result2[output_offset + c_node] & btrofr;

      /* Check all output edges */
      if ((out_edges[c_node].x != 0) && (out_edges[c_node].y != 0)) {
        unsigned int result1 = 0;
        unsigned int result2 = 0;

        unsigned int t = 0;
        /*out_edges[c_node].x; */
        if (edges[out_edges[c_node].x].x == c_node)
          t = robust_result1[output_offset + out_edges[c_node].x];
        else
          t = robust_result2[output_offset + out_edges[c_node].x];
        result1 = t & temp_final_result &
                  is_robust_path_to_po[output_offset + out_edges[c_node].x];

        if (edges[out_edges[c_node].y].x == c_node)
          t = robust_result1[output_offset + out_edges[c_node].y];
        else
          t = robust_result2[output_offset + out_edges[c_node].y];
        result2 = t & temp_final_result &
                  is_robust_path_to_po[output_offset + out_edges[c_node].y];

        p1[output_offset + c_node] = result1;
        p2[output_offset + c_node] = result2;

      } else {
        p1[output_offset + c_node] = temp_final_result;
        p2[output_offset + c_node] = temp_final_result;
      }
    } else {
      p1[output_offset + c_node] = is_robust_path_to_po[output_offset + c_node];
      p2[output_offset + c_node] = is_robust_path_to_po[output_offset + c_node];
    }
  }
}