text: var areaJSON='{"columns":["", "ALUTs", "FFs", "RAMs", "DSPs", "MLABs", "Details"], "debug_enabled":"true", "type":"module", "total_percent":[50.0347, 25.9314, 25.6368, 23.9997, 18.4549], "total":[483302, 956887, 2813, 1063, 32], "name":"Kernel System", "max_resources":[1866240, 3732480, 11721, 5760, 93312], "children":[{"name":"Static Partition", "type":"partition", "children":[{"name":"Board interface", "type":"resource", "data":[455240, 910480, 2627, 1047, 0], "details":[{"type":"text", "text":"Platform interface logic."}]}]}, {"name":"Global interconnect", "type":"resource", "data":[19004, 21208, 0, 0, 0], "details":[{"type":"text", "text":"Global interconnect for 4 global loads and 4 global stores. Reduce number of global loads and stores to simplify global interconnect."}, {"type":"brief", "text":"For 4 global loads and 4 global stores."}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Global Memory Interconnect", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#hnj1476724450050"}]}]}, {"name":"System description ROM", "type":"resource", "data":[2, 71, 2, 0, 0], "details":[{"type":"text", "text":"This read-only memory contains information for the host about the system and about each kernel in the system, including kernel names, arguments, and attributes. The system description ROM ensures that the binary image on the FPGA is compatible with the host program."}, {"type":"brief", "text":"Contains information for the host."}]}, {"name":"accessMemory_0", "compute_units":1, "type":"function", "total_percent":[0.294502, 0.129887, 0.168306, 0.392458, 0.0694444], "total_kernel_resources":[2264, 6282, 46, 4, 8], "details":[{"type":"text", "text":"Number of compute units: 1"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}, {"type":"brief", "text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}], "children":[{"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:47)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:62)", "type":"resource", "data":[40, 128, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":47}], [{"filename":"replicated_random_access_kernels_single.cl", "line":62}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 64"}, {"type":"brief", "text":"Register,\\n1 reg, 64 width"}]}, {"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:58)", "type":"resource", "data":[8, 130, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 65"}, {"type":"brief", "text":"Register,\\n1 reg, 65 width"}]}, {"name":"accessMemory_0.B0", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[4, 704, 0, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[4, 513, 0, 0, 0]}, {"name":"replicated_random_access_kernels_single.cl:51", "type":"resource", "data":[0, 64, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":51}]]}, {"name":"replicated_random_access_kernels_single.cl:58", "type":"resource", "data":[0, 66, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]]}, {"name":"replicated_random_access_kernels_single.cl:73", "type":"resource", "data":[0, 61, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":73}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[10, 9, 1, 0, 0], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:51", "type":"resource", "data":[257, 338, 0, 4, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":51}]], "children":[{"name":"64-bit Integer Multiply", "type":"resource", "count":1, "data":[257, 338, 0, 4, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:58", "type":"resource", "data":[249, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"64-bit Integer Compare", "type":"resource", "count":2, "data":[88, 1, 0, 0, 0]}, {"name":"64-bit Select", "type":"resource", "count":1, "data":[32, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:73", "type":"resource", "data":[61, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":73}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[61, 0, 0, 0, 0]}], "replace_name":"true"}]}]}, {"name":"accessMemory_0.B1", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[154, 530, 9, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[154, 530, 9, 0, 0]}]}, {"name":"Feedback", "type":"resource", "data":[81, 9, 0, 0, 0], "details":[{"type":"brief", "text":"Loop-carried dependencies"}, {"type":"text", "text":"Resources for loop-carried dependencies. To reduce this area:", "details":[{"type":"text", "text":"reduce number and size of loop-carried variables"}]}], "children":[{"name":"replicated_random_access_kernels_single.cl:47", "type":"resource", "data":[65, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":47}]]}, {"name":"replicated_random_access_kernels_single.cl:58", "type":"resource", "data":[16, 9, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[120, 212, 4, 0, 4], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:47", "type":"resource", "data":[52, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":47}]], "children":[{"name":"65-bit Select", "type":"resource", "count":1, "data":[52, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:58", "type":"resource", "data":[88.5, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]], "children":[{"name":"1-bit And", "type":"resource", "count":1, "data":[1, 0, 0, 0, 0]}, {"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}, {"name":"65-bit Integer Compare", "type":"resource", "count":1, "data":[22, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:71", "type":"resource", "data":[63, 63, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":71}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[63, 63, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:73", "type":"resource", "data":[62, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":73}]], "children":[{"name":"64-bit And", "type":"resource", "count":1, "data":[62, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:75", "type":"resource", "data":[64, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":75}]], "children":[{"name":"64-bit Integer Subtract", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:89", "type":"resource", "data":[67, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":89}]], "children":[{"name":"64-bit Integer Compare", "type":"resource", "count":1, "data":[67, 1, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:90", "type":"resource", "data":[453.5, 1985, 15, 0, 4], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":90}]], "children":[{"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"Load", "type":"resource", "count":1, "data":[453, 1985, 15, 0, 4], "details":[{"type":"text", "text":"Load uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:102", "type":"resource", "data":[430, 2172, 17, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":102}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"Store", "type":"resource", "count":1, "data":[366, 2172, 17, 0, 0], "details":[{"type":"text", "text":"Store uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}]}]}]}, {"name":"accessMemory_1", "compute_units":1, "type":"function", "total_percent":[0.294502, 0.129887, 0.168306, 0.392458, 0.0694444], "total_kernel_resources":[2264, 6282, 46, 4, 8], "details":[{"type":"text", "text":"Number of compute units: 1"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}, {"type":"brief", "text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}], "children":[{"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:155)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:170)", "type":"resource", "data":[40, 128, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":155}], [{"filename":"replicated_random_access_kernels_single.cl", "line":170}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 64"}, {"type":"brief", "text":"Register,\\n1 reg, 64 width"}]}, {"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:166)", "type":"resource", "data":[8, 130, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 65"}, {"type":"brief", "text":"Register,\\n1 reg, 65 width"}]}, {"name":"accessMemory_1.B0", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[4, 704, 0, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[4, 513, 0, 0, 0]}, {"name":"replicated_random_access_kernels_single.cl:159", "type":"resource", "data":[0, 64, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":159}]]}, {"name":"replicated_random_access_kernels_single.cl:166", "type":"resource", "data":[0, 66, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]]}, {"name":"replicated_random_access_kernels_single.cl:181", "type":"resource", "data":[0, 61, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":181}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[10, 9, 1, 0, 0], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:159", "type":"resource", "data":[257, 338, 0, 4, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":159}]], "children":[{"name":"64-bit Integer Multiply", "type":"resource", "count":1, "data":[257, 338, 0, 4, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:166", "type":"resource", "data":[249, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"64-bit Integer Compare", "type":"resource", "count":2, "data":[88, 1, 0, 0, 0]}, {"name":"64-bit Select", "type":"resource", "count":1, "data":[32, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:181", "type":"resource", "data":[61, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":181}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[61, 0, 0, 0, 0]}], "replace_name":"true"}]}]}, {"name":"accessMemory_1.B1", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[154, 530, 9, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[154, 530, 9, 0, 0]}]}, {"name":"Feedback", "type":"resource", "data":[81, 9, 0, 0, 0], "details":[{"type":"brief", "text":"Loop-carried dependencies"}, {"type":"text", "text":"Resources for loop-carried dependencies. To reduce this area:", "details":[{"type":"text", "text":"reduce number and size of loop-carried variables"}]}], "children":[{"name":"replicated_random_access_kernels_single.cl:155", "type":"resource", "data":[65, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":155}]]}, {"name":"replicated_random_access_kernels_single.cl:166", "type":"resource", "data":[16, 9, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[120, 212, 4, 0, 4], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:155", "type":"resource", "data":[52, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":155}]], "children":[{"name":"65-bit Select", "type":"resource", "count":1, "data":[52, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:166", "type":"resource", "data":[88.5, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]], "children":[{"name":"1-bit And", "type":"resource", "count":1, "data":[1, 0, 0, 0, 0]}, {"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}, {"name":"65-bit Integer Compare", "type":"resource", "count":1, "data":[22, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:179", "type":"resource", "data":[63, 63, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":179}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[63, 63, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:181", "type":"resource", "data":[62, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":181}]], "children":[{"name":"64-bit And", "type":"resource", "count":1, "data":[62, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:183", "type":"resource", "data":[64, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":183}]], "children":[{"name":"64-bit Integer Subtract", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:197", "type":"resource", "data":[67, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":197}]], "children":[{"name":"64-bit Integer Compare", "type":"resource", "count":1, "data":[67, 1, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:198", "type":"resource", "data":[453.5, 1985, 15, 0, 4], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":198}]], "children":[{"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"Load", "type":"resource", "count":1, "data":[453, 1985, 15, 0, 4], "details":[{"type":"text", "text":"Load uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:210", "type":"resource", "data":[430, 2172, 17, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":210}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"Store", "type":"resource", "count":1, "data":[366, 2172, 17, 0, 0], "details":[{"type":"text", "text":"Store uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}]}]}]}, {"name":"accessMemory_2", "compute_units":1, "type":"function", "total_percent":[0.294502, 0.129887, 0.168306, 0.392458, 0.0694444], "total_kernel_resources":[2264, 6282, 46, 4, 8], "details":[{"type":"text", "text":"Number of compute units: 1"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}, {"type":"brief", "text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}], "children":[{"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:263)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:278)", "type":"resource", "data":[40, 128, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":263}], [{"filename":"replicated_random_access_kernels_single.cl", "line":278}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 64"}, {"type":"brief", "text":"Register,\\n1 reg, 64 width"}]}, {"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:274)", "type":"resource", "data":[8, 130, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 65"}, {"type":"brief", "text":"Register,\\n1 reg, 65 width"}]}, {"name":"accessMemory_2.B0", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[4, 704, 0, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[4, 513, 0, 0, 0]}, {"name":"replicated_random_access_kernels_single.cl:267", "type":"resource", "data":[0, 64, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":267}]]}, {"name":"replicated_random_access_kernels_single.cl:274", "type":"resource", "data":[0, 66, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]]}, {"name":"replicated_random_access_kernels_single.cl:289", "type":"resource", "data":[0, 61, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":289}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[10, 9, 1, 0, 0], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:267", "type":"resource", "data":[257, 338, 0, 4, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":267}]], "children":[{"name":"64-bit Integer Multiply", "type":"resource", "count":1, "data":[257, 338, 0, 4, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:274", "type":"resource", "data":[249, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"64-bit Integer Compare", "type":"resource", "count":2, "data":[88, 1, 0, 0, 0]}, {"name":"64-bit Select", "type":"resource", "count":1, "data":[32, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:289", "type":"resource", "data":[61, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":289}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[61, 0, 0, 0, 0]}], "replace_name":"true"}]}]}, {"name":"accessMemory_2.B1", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[154, 530, 9, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[154, 530, 9, 0, 0]}]}, {"name":"Feedback", "type":"resource", "data":[81, 9, 0, 0, 0], "details":[{"type":"brief", "text":"Loop-carried dependencies"}, {"type":"text", "text":"Resources for loop-carried dependencies. To reduce this area:", "details":[{"type":"text", "text":"reduce number and size of loop-carried variables"}]}], "children":[{"name":"replicated_random_access_kernels_single.cl:263", "type":"resource", "data":[65, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":263}]]}, {"name":"replicated_random_access_kernels_single.cl:274", "type":"resource", "data":[16, 9, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[120, 212, 4, 0, 4], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:263", "type":"resource", "data":[52, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":263}]], "children":[{"name":"65-bit Select", "type":"resource", "count":1, "data":[52, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:274", "type":"resource", "data":[88.5, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]], "children":[{"name":"1-bit And", "type":"resource", "count":1, "data":[1, 0, 0, 0, 0]}, {"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}, {"name":"65-bit Integer Compare", "type":"resource", "count":1, "data":[22, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:287", "type":"resource", "data":[63, 63, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":287}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[63, 63, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:289", "type":"resource", "data":[62, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":289}]], "children":[{"name":"64-bit And", "type":"resource", "count":1, "data":[62, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:291", "type":"resource", "data":[64, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":291}]], "children":[{"name":"64-bit Integer Subtract", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:305", "type":"resource", "data":[67, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":305}]], "children":[{"name":"64-bit Integer Compare", "type":"resource", "count":1, "data":[67, 1, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:306", "type":"resource", "data":[453.5, 1985, 15, 0, 4], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":306}]], "children":[{"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"Load", "type":"resource", "count":1, "data":[453, 1985, 15, 0, 4], "details":[{"type":"text", "text":"Load uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:318", "type":"resource", "data":[430, 2172, 17, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":318}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"Store", "type":"resource", "count":1, "data":[366, 2172, 17, 0, 0], "details":[{"type":"text", "text":"Store uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}]}]}]}, {"name":"accessMemory_3", "compute_units":1, "type":"function", "total_percent":[0.294502, 0.129887, 0.168306, 0.392458, 0.0694444], "total_kernel_resources":[2264, 6282, 46, 4, 8], "details":[{"type":"text", "text":"Number of compute units: 1"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}, {"type":"brief", "text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality."}], "children":[{"name":"Coalesced Private Variables: \\n - \'update_val\' (replicated_random_access_kernels_single.cl:386)\\n - \'ran\' (replicated_random_access_kernels_single.cl:371)", "type":"resource", "data":[40, 128, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":371}], [{"filename":"replicated_random_access_kernels_single.cl", "line":386}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 64"}, {"type":"brief", "text":"Register,\\n1 reg, 64 width"}]}, {"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:382)", "type":"resource", "data":[8, 130, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]], "details":[{"type":"text", "text":"Type: Register"}, {"type":"text", "text":"1 register of width 65"}, {"type":"brief", "text":"Register,\\n1 reg, 65 width"}]}, {"name":"accessMemory_3.B0", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[4, 704, 0, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[4, 513, 0, 0, 0]}, {"name":"replicated_random_access_kernels_single.cl:375", "type":"resource", "data":[0, 64, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":375}]]}, {"name":"replicated_random_access_kernels_single.cl:382", "type":"resource", "data":[0, 66, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]]}, {"name":"replicated_random_access_kernels_single.cl:397", "type":"resource", "data":[0, 61, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":397}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[10, 9, 1, 0, 0], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:375", "type":"resource", "data":[257, 338, 0, 4, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":375}]], "children":[{"name":"64-bit Integer Multiply", "type":"resource", "count":1, "data":[257, 338, 0, 4, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:382", "type":"resource", "data":[249, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"64-bit Integer Compare", "type":"resource", "count":2, "data":[88, 1, 0, 0, 0]}, {"name":"64-bit Select", "type":"resource", "count":1, "data":[32, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:397", "type":"resource", "data":[61, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":397}]], "children":[{"name":"64-bit Integer Add", "type":"resource", "count":1, "data":[61, 0, 0, 0, 0]}], "replace_name":"true"}]}]}, {"name":"accessMemory_3.B1", "type":"basicblock", "children":[{"name":"State", "type":"resource", "data":[154, 530, 9, 0, 0], "details":[{"type":"brief", "text":"Live values and control logic"}, {"type":"text", "text":"Resources for live values and control logic. To reduce this area:", "details":[{"type":"text", "text":"reduce size of local variables"}, {"type":"text", "text":"reduce scope of local variables, localizing them as much as possible"}, {"type":"text", "text":"reduce number of nested loops"}]}], "children":[{"name":"No Source Line", "type":"resource", "data":[154, 530, 9, 0, 0]}]}, {"name":"Feedback", "type":"resource", "data":[81, 9, 0, 0, 0], "details":[{"type":"brief", "text":"Loop-carried dependencies"}, {"type":"text", "text":"Resources for loop-carried dependencies. To reduce this area:", "details":[{"type":"text", "text":"reduce number and size of loop-carried variables"}]}], "children":[{"name":"replicated_random_access_kernels_single.cl:371", "type":"resource", "data":[65, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":371}]]}, {"name":"replicated_random_access_kernels_single.cl:382", "type":"resource", "data":[16, 9, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]]}]}, {"name":"Cluster logic", "type":"resource", "data":[120, 212, 4, 0, 4], "details":[{"type":"brief", "text":"Logic required to efficiently support sets of operations that do not stall"}, {"type":"text", "text":"Logic required to efficiently support sets of operations that do not stall. This area cannot be affected directly."}]}, {"name":"Computation", "type":"resource", "children":[{"name":"replicated_random_access_kernels_single.cl:371", "type":"resource", "data":[52, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":371}]], "children":[{"name":"65-bit Select", "type":"resource", "count":1, "data":[52, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:382", "type":"resource", "data":[88.5, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]], "children":[{"name":"1-bit And", "type":"resource", "count":1, "data":[1, 0, 0, 0, 0]}, {"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"65-bit Integer Add", "type":"resource", "count":1, "data":[65, 0, 0, 0, 0]}, {"name":"65-bit Integer Compare", "type":"resource", "count":1, "data":[22, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:395", "type":"resource", "data":[63, 63, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":395}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[63, 63, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:397", "type":"resource", "data":[62, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":397}]], "children":[{"name":"64-bit And", "type":"resource", "count":1, "data":[62, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:399", "type":"resource", "data":[64, 0, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":399}]], "children":[{"name":"64-bit Integer Subtract", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:413", "type":"resource", "data":[67, 1, 0, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":413}]], "children":[{"name":"64-bit Integer Compare", "type":"resource", "count":1, "data":[67, 1, 0, 0, 0]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:414", "type":"resource", "data":[453.5, 1985, 15, 0, 4], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":414}]], "children":[{"name":"1-bit Or", "type":"resource", "count":1, "data":[0.5, 0, 0, 0, 0]}, {"name":"Load", "type":"resource", "count":1, "data":[453, 1985, 15, 0, 4], "details":[{"type":"text", "text":"Load uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}, {"name":"replicated_random_access_kernels_single.cl:426", "type":"resource", "data":[430, 2172, 17, 0, 0], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":426}]], "children":[{"name":"64-bit Xor", "type":"resource", "count":1, "data":[64, 0, 0, 0, 0]}, {"name":"Store", "type":"resource", "count":1, "data":[366, 2172, 17, 0, 0], "details":[{"type":"text", "text":"Store uses a Burst-coalesced LSU"}, {"type":"brief", "text":"Burst-coalesced LSU"}]}], "replace_name":"true"}]}]}]}]}';
  var area_srcJSON='{"children":[{"children":[{"data":[455240,910480,2627,1047,0],"details":[{"text":"Platform interface logic.","type":"text"}],"name":"Board interface","type":"resource"}],"name":"Static Partition","type":"partition"},{"data":[19004,21208,0,0,0],"details":[{"text":"Global interconnect for 4 global loads and 4 global stores. Reduce number of global loads and stores to simplify global interconnect.","type":"text"},{"text":"For 4 global loads and 4 global stores.","type":"brief"},{"links":[{"guide":"Best Practices Guide : Global Memory Interconnect","link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#hnj1476724450050"}],"text":"See %L for more information","type":"text"}],"name":"Global interconnect","type":"resource"},{"data":[2,71,2,0,0],"details":[{"text":"This read-only memory contains information for the host about the system and about each kernel in the system, including kernel names, arguments, and attributes. The system description ROM ensures that the binary image on the FPGA is compatible with the host program.","type":"text"},{"text":"Contains information for the host.","type":"brief"}],"name":"System description ROM","type":"resource"},{"children":[{"data":[211,230,5,0,4],"details":[{"text":"Feedback+Cluster logic","type":"brief"}],"name":"Data control overhead","type":"resource"},{"data":[40,128,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 64","type":"text"},{"text":"Register,\\n1 reg, 64 width","type":"brief"}],"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:47)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:62)","type":"resource"},{"data":[8,130,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 65","type":"text"},{"text":"Register,\\n1 reg, 65 width","type":"brief"}],"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:58)","type":"resource"},{"children":[{"count":2,"data":[158,1043,9,0,0],"debug":[[{"filename":"","line":0}]],"name":"State","type":"resource"}],"data":[158,1043,9,0,0],"name":"No Source Line","type":"resource"},{"children":[{"count":1,"data":[0,64,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":51}]],"name":"State","type":"resource"},{"count":1,"data":[257,338,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":51}]],"name":"64-bit Integer Multiply","type":"resource"}],"data":[257,402,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":51}]],"name":"replicated_random_access_kernels_single.cl:51","type":"resource"},{"children":[{"count":1,"data":[0,66,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"State","type":"resource"},{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"64-bit Integer Add","type":"resource"},{"count":2,"data":[88,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"64-bit Integer Compare","type":"resource"},{"count":1,"data":[32,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"64-bit Select","type":"resource"},{"count":2,"data":[130,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"65-bit Integer Add","type":"resource"},{"count":1,"data":[1,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"1-bit And","type":"resource"},{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[22,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"65-bit Integer Compare","type":"resource"}],"data":[337.5,67,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":58}]],"name":"replicated_random_access_kernels_single.cl:58","type":"resource"},{"children":[{"count":1,"data":[0,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":73}]],"name":"State","type":"resource"},{"count":1,"data":[61,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":73}]],"name":"64-bit Integer Add","type":"resource"},{"count":1,"data":[62,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":73}]],"name":"64-bit And","type":"resource"}],"data":[123,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":73}]],"name":"replicated_random_access_kernels_single.cl:73","type":"resource"},{"children":[{"count":1,"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":47}]],"name":"65-bit Select","type":"resource"}],"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":47}]],"name":"replicated_random_access_kernels_single.cl:47","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":71}]],"name":"64-bit Xor","type":"resource"}],"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":71}]],"name":"replicated_random_access_kernels_single.cl:71","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":75}]],"name":"64-bit Integer Subtract","type":"resource"}],"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":75}]],"name":"replicated_random_access_kernels_single.cl:75","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":89}]],"name":"64-bit Integer Compare","type":"resource"}],"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":89}]],"name":"replicated_random_access_kernels_single.cl:89","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":90}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[453,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":90}]],"name":"Load","type":"resource"}],"data":[453.5,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":90}]],"name":"replicated_random_access_kernels_single.cl:90","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":102}]],"name":"64-bit Xor","type":"resource"},{"count":1,"data":[366,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":102}]],"name":"Store","type":"resource"}],"data":[430,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":102}]],"name":"replicated_random_access_kernels_single.cl:102","replace_name":"true","type":"resource"}],"compute_units":1,"data":[2264,6282,46,4,8],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":47}]],"details":[{"text":"Number of compute units: 1","type":"text"},{"text":"Max global work dimension: 0","type":"text"},{"text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"text"},{"text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"brief"}],"name":"accessMemory_0","total_kernel_resources":[2264,6282,46,4,8],"total_percent":[0.294502,0.129887,0.168306,0.392458,0.0694444],"type":"function"},{"children":[{"data":[211,230,5,0,4],"details":[{"text":"Feedback+Cluster logic","type":"brief"}],"name":"Data control overhead","type":"resource"},{"data":[40,128,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 64","type":"text"},{"text":"Register,\\n1 reg, 64 width","type":"brief"}],"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:155)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:170)","type":"resource"},{"data":[8,130,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 65","type":"text"},{"text":"Register,\\n1 reg, 65 width","type":"brief"}],"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:166)","type":"resource"},{"children":[{"count":2,"data":[158,1043,9,0,0],"debug":[[{"filename":"","line":0}]],"name":"State","type":"resource"}],"data":[158,1043,9,0,0],"name":"No Source Line","type":"resource"},{"children":[{"count":1,"data":[0,64,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":159}]],"name":"State","type":"resource"},{"count":1,"data":[257,338,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":159}]],"name":"64-bit Integer Multiply","type":"resource"}],"data":[257,402,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":159}]],"name":"replicated_random_access_kernels_single.cl:159","type":"resource"},{"children":[{"count":1,"data":[0,66,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"State","type":"resource"},{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"64-bit Integer Add","type":"resource"},{"count":2,"data":[88,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"64-bit Integer Compare","type":"resource"},{"count":1,"data":[32,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"64-bit Select","type":"resource"},{"count":2,"data":[130,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"65-bit Integer Add","type":"resource"},{"count":1,"data":[1,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"1-bit And","type":"resource"},{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[22,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"65-bit Integer Compare","type":"resource"}],"data":[337.5,67,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":166}]],"name":"replicated_random_access_kernels_single.cl:166","type":"resource"},{"children":[{"count":1,"data":[0,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":181}]],"name":"State","type":"resource"},{"count":1,"data":[61,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":181}]],"name":"64-bit Integer Add","type":"resource"},{"count":1,"data":[62,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":181}]],"name":"64-bit And","type":"resource"}],"data":[123,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":181}]],"name":"replicated_random_access_kernels_single.cl:181","type":"resource"},{"children":[{"count":1,"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":155}]],"name":"65-bit Select","type":"resource"}],"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":155}]],"name":"replicated_random_access_kernels_single.cl:155","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":179}]],"name":"64-bit Xor","type":"resource"}],"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":179}]],"name":"replicated_random_access_kernels_single.cl:179","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":183}]],"name":"64-bit Integer Subtract","type":"resource"}],"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":183}]],"name":"replicated_random_access_kernels_single.cl:183","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":197}]],"name":"64-bit Integer Compare","type":"resource"}],"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":197}]],"name":"replicated_random_access_kernels_single.cl:197","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":198}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[453,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":198}]],"name":"Load","type":"resource"}],"data":[453.5,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":198}]],"name":"replicated_random_access_kernels_single.cl:198","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":210}]],"name":"64-bit Xor","type":"resource"},{"count":1,"data":[366,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":210}]],"name":"Store","type":"resource"}],"data":[430,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":210}]],"name":"replicated_random_access_kernels_single.cl:210","replace_name":"true","type":"resource"}],"compute_units":1,"data":[2264,6282,46,4,8],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":155}]],"details":[{"text":"Number of compute units: 1","type":"text"},{"text":"Max global work dimension: 0","type":"text"},{"text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"text"},{"text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"brief"}],"name":"accessMemory_1","total_kernel_resources":[2264,6282,46,4,8],"total_percent":[0.294502,0.129887,0.168306,0.392458,0.0694444],"type":"function"},{"children":[{"data":[211,230,5,0,4],"details":[{"text":"Feedback+Cluster logic","type":"brief"}],"name":"Data control overhead","type":"resource"},{"data":[40,128,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 64","type":"text"},{"text":"Register,\\n1 reg, 64 width","type":"brief"}],"name":"Coalesced Private Variables: \\n - \'ran\' (replicated_random_access_kernels_single.cl:263)\\n - \'update_val\' (replicated_random_access_kernels_single.cl:278)","type":"resource"},{"data":[8,130,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 65","type":"text"},{"text":"Register,\\n1 reg, 65 width","type":"brief"}],"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:274)","type":"resource"},{"children":[{"count":2,"data":[158,1043,9,0,0],"debug":[[{"filename":"","line":0}]],"name":"State","type":"resource"}],"data":[158,1043,9,0,0],"name":"No Source Line","type":"resource"},{"children":[{"count":1,"data":[0,64,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":267}]],"name":"State","type":"resource"},{"count":1,"data":[257,338,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":267}]],"name":"64-bit Integer Multiply","type":"resource"}],"data":[257,402,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":267}]],"name":"replicated_random_access_kernels_single.cl:267","type":"resource"},{"children":[{"count":1,"data":[0,66,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"State","type":"resource"},{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"64-bit Integer Add","type":"resource"},{"count":2,"data":[88,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"64-bit Integer Compare","type":"resource"},{"count":1,"data":[32,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"64-bit Select","type":"resource"},{"count":2,"data":[130,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"65-bit Integer Add","type":"resource"},{"count":1,"data":[1,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"1-bit And","type":"resource"},{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[22,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"65-bit Integer Compare","type":"resource"}],"data":[337.5,67,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":274}]],"name":"replicated_random_access_kernels_single.cl:274","type":"resource"},{"children":[{"count":1,"data":[0,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":289}]],"name":"State","type":"resource"},{"count":1,"data":[61,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":289}]],"name":"64-bit Integer Add","type":"resource"},{"count":1,"data":[62,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":289}]],"name":"64-bit And","type":"resource"}],"data":[123,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":289}]],"name":"replicated_random_access_kernels_single.cl:289","type":"resource"},{"children":[{"count":1,"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":263}]],"name":"65-bit Select","type":"resource"}],"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":263}]],"name":"replicated_random_access_kernels_single.cl:263","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":287}]],"name":"64-bit Xor","type":"resource"}],"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":287}]],"name":"replicated_random_access_kernels_single.cl:287","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":291}]],"name":"64-bit Integer Subtract","type":"resource"}],"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":291}]],"name":"replicated_random_access_kernels_single.cl:291","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":305}]],"name":"64-bit Integer Compare","type":"resource"}],"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":305}]],"name":"replicated_random_access_kernels_single.cl:305","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":306}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[453,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":306}]],"name":"Load","type":"resource"}],"data":[453.5,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":306}]],"name":"replicated_random_access_kernels_single.cl:306","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":318}]],"name":"64-bit Xor","type":"resource"},{"count":1,"data":[366,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":318}]],"name":"Store","type":"resource"}],"data":[430,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":318}]],"name":"replicated_random_access_kernels_single.cl:318","replace_name":"true","type":"resource"}],"compute_units":1,"data":[2264,6282,46,4,8],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":263}]],"details":[{"text":"Number of compute units: 1","type":"text"},{"text":"Max global work dimension: 0","type":"text"},{"text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"text"},{"text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"brief"}],"name":"accessMemory_2","total_kernel_resources":[2264,6282,46,4,8],"total_percent":[0.294502,0.129887,0.168306,0.392458,0.0694444],"type":"function"},{"children":[{"data":[211,230,5,0,4],"details":[{"text":"Feedback+Cluster logic","type":"brief"}],"name":"Data control overhead","type":"resource"},{"data":[40,128,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 64","type":"text"},{"text":"Register,\\n1 reg, 64 width","type":"brief"}],"name":"Coalesced Private Variables: \\n - \'update_val\' (replicated_random_access_kernels_single.cl:386)\\n - \'ran\' (replicated_random_access_kernels_single.cl:371)","type":"resource"},{"data":[8,130,0,0,0],"details":[{"text":"Type: Register","type":"text"},{"text":"1 register of width 65","type":"text"},{"text":"Register,\\n1 reg, 65 width","type":"brief"}],"name":"Private Variable: \\n - \'i\' (replicated_random_access_kernels_single.cl:382)","type":"resource"},{"children":[{"count":2,"data":[158,1043,9,0,0],"debug":[[{"filename":"","line":0}]],"name":"State","type":"resource"}],"data":[158,1043,9,0,0],"name":"No Source Line","type":"resource"},{"children":[{"count":1,"data":[0,64,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":375}]],"name":"State","type":"resource"},{"count":1,"data":[257,338,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":375}]],"name":"64-bit Integer Multiply","type":"resource"}],"data":[257,402,0,4,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":375}]],"name":"replicated_random_access_kernels_single.cl:375","type":"resource"},{"children":[{"count":1,"data":[0,66,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"State","type":"resource"},{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"64-bit Integer Add","type":"resource"},{"count":2,"data":[88,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"64-bit Integer Compare","type":"resource"},{"count":1,"data":[32,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"64-bit Select","type":"resource"},{"count":2,"data":[130,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"65-bit Integer Add","type":"resource"},{"count":1,"data":[1,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"1-bit And","type":"resource"},{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[22,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"65-bit Integer Compare","type":"resource"}],"data":[337.5,67,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":382}]],"name":"replicated_random_access_kernels_single.cl:382","type":"resource"},{"children":[{"count":1,"data":[0,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":397}]],"name":"State","type":"resource"},{"count":1,"data":[61,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":397}]],"name":"64-bit Integer Add","type":"resource"},{"count":1,"data":[62,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":397}]],"name":"64-bit And","type":"resource"}],"data":[123,61,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":397}]],"name":"replicated_random_access_kernels_single.cl:397","type":"resource"},{"children":[{"count":1,"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":371}]],"name":"65-bit Select","type":"resource"}],"data":[52,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":371}]],"name":"replicated_random_access_kernels_single.cl:371","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":395}]],"name":"64-bit Xor","type":"resource"}],"data":[63,63,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":395}]],"name":"replicated_random_access_kernels_single.cl:395","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":399}]],"name":"64-bit Integer Subtract","type":"resource"}],"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":399}]],"name":"replicated_random_access_kernels_single.cl:399","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":413}]],"name":"64-bit Integer Compare","type":"resource"}],"data":[67,1,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":413}]],"name":"replicated_random_access_kernels_single.cl:413","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[0.5,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":414}]],"name":"1-bit Or","type":"resource"},{"count":1,"data":[453,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":414}]],"name":"Load","type":"resource"}],"data":[453.5,1985,15,0,4],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":414}]],"name":"replicated_random_access_kernels_single.cl:414","replace_name":"true","type":"resource"},{"children":[{"count":1,"data":[64,0,0,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":426}]],"name":"64-bit Xor","type":"resource"},{"count":1,"data":[366,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":426}]],"name":"Store","type":"resource"}],"data":[430,2172,17,0,0],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":426}]],"name":"replicated_random_access_kernels_single.cl:426","replace_name":"true","type":"resource"}],"compute_units":1,"data":[2264,6282,46,4,8],"debug":[[{"filename":"replicated_random_access_kernels_single.cl","line":371}]],"details":[{"text":"Number of compute units: 1","type":"text"},{"text":"Max global work dimension: 0","type":"text"},{"text":"Kernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"text"},{"text":"1 compute unit.\\nMax global work dimension: 0\\nKernel attribute \'uses_global_work_offset\' not specified. Add \'__attribute__((uses_global_work_offset(0)))\' to reduce area utilization if the kernel will always be enqueued with a 0 or NULL global_work_offset argument in clEnqueueNDRangeKernel. Add \'__attribute__((uses_global_work_offset(1)))\' to hide this suggestion without affecting kernel functionality.","type":"brief"}],"name":"accessMemory_3","total_kernel_resources":[2264,6282,46,4,8],"total_percent":[0.294502,0.129887,0.168306,0.392458,0.0694444],"type":"function"}],"columns":["","ALUTs","FFs","RAMs","DSPs","MLABs","Details"],"data":[28062,46407,186,16,32],"debug_enabled":"true","max_resources":[1866240,3732480,11721,5760,93312],"name":"Kernel System","total":[483302,956887,2813,1063,32],"total_percent":[50.0347,25.9314,25.6368,23.9997,18.4549],"type":"module"}';
  var mavJSON='{"nodes":[{"type":"kernel", "id":2, "name":"accessMemory_0", "children":[{"type":"bb", "id":3, "name":"accessMemory_0.B0", "details":[{"type":"table", "Latency":"42"}]}, {"type":"bb", "id":4, "name":"accessMemory_0.B1", "children":[{"type":"inst", "id":6, "name":"Load", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":90}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"40", "Latency":"512", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":7, "name":"Store", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":102}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"583", "Latency":"64", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":8, "name":"Loop Input", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]], "details":[{"type":"table", "Start Cycle":"0", "Latency":"1", "Loops To":"9"}]}, {"type":"inst", "id":9, "name":"Loop End", "details":[{"type":"table", "Start Cycle":"647", "Latency":"1"}]}], "details":[{"type":"table", "Latency":"647", "II":"1", "Subloops":"No", "Pipelined":"Yes", "Fmax Bottlenecks":"No", "Loop Info":""}]}, {"type":"bb", "id":5, "name":"accessMemory_0.B2", "details":[{"type":"table", "Latency":"0"}]}]}, {"type":"kernel", "id":11, "name":"accessMemory_1", "children":[{"type":"bb", "id":12, "name":"accessMemory_1.B0", "details":[{"type":"table", "Latency":"42"}]}, {"type":"bb", "id":13, "name":"accessMemory_1.B1", "children":[{"type":"inst", "id":15, "name":"Load", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":198}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"40", "Latency":"512", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":16, "name":"Store", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":210}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"583", "Latency":"64", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":17, "name":"Loop Input", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]], "details":[{"type":"table", "Start Cycle":"0", "Latency":"1", "Loops To":"18"}]}, {"type":"inst", "id":18, "name":"Loop End", "details":[{"type":"table", "Start Cycle":"647", "Latency":"1"}]}], "details":[{"type":"table", "Latency":"647", "II":"1", "Subloops":"No", "Pipelined":"Yes", "Fmax Bottlenecks":"No", "Loop Info":""}]}, {"type":"bb", "id":14, "name":"accessMemory_1.B2", "details":[{"type":"table", "Latency":"0"}]}]}, {"type":"kernel", "id":19, "name":"accessMemory_2", "children":[{"type":"bb", "id":20, "name":"accessMemory_2.B0", "details":[{"type":"table", "Latency":"42"}]}, {"type":"bb", "id":21, "name":"accessMemory_2.B1", "children":[{"type":"inst", "id":23, "name":"Load", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":306}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"40", "Latency":"512", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":24, "name":"Store", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":318}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"583", "Latency":"64", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":25, "name":"Loop Input", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]], "details":[{"type":"table", "Start Cycle":"0", "Latency":"1", "Loops To":"26"}]}, {"type":"inst", "id":26, "name":"Loop End", "details":[{"type":"table", "Start Cycle":"647", "Latency":"1"}]}], "details":[{"type":"table", "Latency":"647", "II":"1", "Subloops":"No", "Pipelined":"Yes", "Fmax Bottlenecks":"No", "Loop Info":""}]}, {"type":"bb", "id":22, "name":"accessMemory_2.B2", "details":[{"type":"table", "Latency":"0"}]}]}, {"type":"kernel", "id":27, "name":"accessMemory_3", "children":[{"type":"bb", "id":28, "name":"accessMemory_3.B0", "details":[{"type":"table", "Latency":"42"}]}, {"type":"bb", "id":29, "name":"accessMemory_3.B1", "children":[{"type":"inst", "id":31, "name":"Load", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":414}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"40", "Latency":"512", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":32, "name":"Store", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":426}]], "details":[{"type":"table", "Width":"64 bits", "Type":"Burst-coalesced", "Stall-free":"No", "Start Cycle":"583", "Latency":"64", "Reference":[{"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Load-Store Units", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#yeo1491314105959"}]}]}]}, {"type":"inst", "id":33, "name":"Loop Input", "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]], "details":[{"type":"table", "Start Cycle":"0", "Latency":"1", "Loops To":"34"}]}, {"type":"inst", "id":34, "name":"Loop End", "details":[{"type":"table", "Start Cycle":"647", "Latency":"1"}]}], "details":[{"type":"table", "Latency":"647", "II":"1", "Subloops":"No", "Pipelined":"Yes", "Fmax Bottlenecks":"No", "Loop Info":""}]}, {"type":"bb", "id":30, "name":"accessMemory_3.B2", "details":[{"type":"table", "Latency":"0"}]}]}, {"type":"memtype", "id":1, "name":"Global Memory", "children":[{"type":"memsys", "id":10, "name":"DDR", "details":[{"type":"table", "Number of banks":"4"}]}]}], "links":[{"from":3, "to":8}, {"from":6, "to":9}, {"from":7, "to":9}, {"from":9, "to":5}, {"from":8, "to":6}, {"from":6, "to":7}, {"from":7, "to":10}, {"from":10, "to":6}, {"from":12, "to":17}, {"from":15, "to":18}, {"from":16, "to":18}, {"from":18, "to":14}, {"from":17, "to":15}, {"from":15, "to":16}, {"from":10, "to":15}, {"from":16, "to":10}, {"from":20, "to":25}, {"from":23, "to":26}, {"from":24, "to":26}, {"from":26, "to":22}, {"from":25, "to":23}, {"from":23, "to":24}, {"from":24, "to":10}, {"from":10, "to":23}, {"from":28, "to":33}, {"from":31, "to":34}, {"from":32, "to":34}, {"from":34, "to":30}, {"from":33, "to":31}, {"from":31, "to":32}, {"from":32, "to":10}, {"from":10, "to":31}]}';
  var loopsJSON='{"columns":["", "Pipelined", "II", "Speculated iterations", "Details"], "children":[{"name":"Kernel: accessMemory_0", "data":["", "", ""], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":42}]], "details":[{"type":"brief", "text":"Single work-item execution"}, {"type":"text", "text":"Single work-item execution"}, {"type":"text", "text":"Fmax bottlenck block: accessMemory_0.B1"}, {"type":"text", "text":"Use the %L report to view more details", "links":[{"view":"Fmax II Report"}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Kernels", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#ipp1476408832230"}]}], "children":[{"name":"accessMemory_0.B1", "data":["Yes", "~1", "9"], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":58}]], "details":[{"type":"brief", "text":" "}, {"type":"text", "text":"Hyper-Optimized loop structure: enabled."}, {"type":"text", "text":"II is an approximation due to the following stallable instructions:", "details":[{"type":"text", "text":"Load Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"90"}]}, {"type":"text", "text":"Store Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"102"}]}]}, {"type":"text", "text":"Maximum concurrent iterations: Capacity of loop", "details":[{"type":"text", "text":"Use the %L viewer to estimate capacity", "links":[{"view":"Fmax II Report"}]}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Loops in a Single Work-Item Kernel", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#wfa1476380079940"}]}], "children":[]}]}, {"name":"Kernel: accessMemory_1", "data":["", "", ""], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":150}]], "details":[{"type":"brief", "text":"Single work-item execution"}, {"type":"text", "text":"Single work-item execution"}, {"type":"text", "text":"Fmax bottlenck block: accessMemory_1.B1"}, {"type":"text", "text":"Use the %L report to view more details", "links":[{"view":"Fmax II Report"}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Kernels", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#ipp1476408832230"}]}], "children":[{"name":"accessMemory_1.B1", "data":["Yes", "~1", "9"], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":166}]], "details":[{"type":"brief", "text":" "}, {"type":"text", "text":"Hyper-Optimized loop structure: enabled."}, {"type":"text", "text":"II is an approximation due to the following stallable instructions:", "details":[{"type":"text", "text":"Load Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"198"}]}, {"type":"text", "text":"Store Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"210"}]}]}, {"type":"text", "text":"Maximum concurrent iterations: Capacity of loop", "details":[{"type":"text", "text":"Use the %L viewer to estimate capacity", "links":[{"view":"Fmax II Report"}]}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Loops in a Single Work-Item Kernel", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#wfa1476380079940"}]}], "children":[]}]}, {"name":"Kernel: accessMemory_2", "data":["", "", ""], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":258}]], "details":[{"type":"brief", "text":"Single work-item execution"}, {"type":"text", "text":"Single work-item execution"}, {"type":"text", "text":"Fmax bottlenck block: accessMemory_2.B1"}, {"type":"text", "text":"Use the %L report to view more details", "links":[{"view":"Fmax II Report"}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Kernels", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#ipp1476408832230"}]}], "children":[{"name":"accessMemory_2.B1", "data":["Yes", "~1", "9"], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":274}]], "details":[{"type":"brief", "text":" "}, {"type":"text", "text":"Hyper-Optimized loop structure: enabled."}, {"type":"text", "text":"II is an approximation due to the following stallable instructions:", "details":[{"type":"text", "text":"Load Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"306"}]}, {"type":"text", "text":"Store Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"318"}]}]}, {"type":"text", "text":"Maximum concurrent iterations: Capacity of loop", "details":[{"type":"text", "text":"Use the %L viewer to estimate capacity", "links":[{"view":"Fmax II Report"}]}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Loops in a Single Work-Item Kernel", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#wfa1476380079940"}]}], "children":[]}]}, {"name":"Kernel: accessMemory_3", "data":["", "", ""], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":366}]], "details":[{"type":"brief", "text":"Single work-item execution"}, {"type":"text", "text":"Single work-item execution"}, {"type":"text", "text":"Fmax bottlenck block: accessMemory_3.B1"}, {"type":"text", "text":"Use the %L report to view more details", "links":[{"view":"Fmax II Report"}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Kernels", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#ipp1476408832230"}]}], "children":[{"name":"accessMemory_3.B1", "data":["Yes", "~1", "9"], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":382}]], "details":[{"type":"brief", "text":" "}, {"type":"text", "text":"Hyper-Optimized loop structure: enabled."}, {"type":"text", "text":"II is an approximation due to the following stallable instructions:", "details":[{"type":"text", "text":"Load Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"414"}]}, {"type":"text", "text":"Store Operation (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"426"}]}]}, {"type":"text", "text":"Maximum concurrent iterations: Capacity of loop", "details":[{"type":"text", "text":"Use the %L viewer to estimate capacity", "links":[{"view":"Fmax II Report"}]}]}, {"type":"text", "text":"See %L for more information", "links":[{"guide":"Best Practices Guide : Loops in a Single Work-Item Kernel", "link":"https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807516407.html#wfa1476380079940"}]}], "children":[]}]}]}';
  var fmax_iiJSON='{"basicblocks":{"accessMemory_0.B0":{"name":"accessMemory_0.B0", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":42, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_0.B1":{"name":"accessMemory_0.B1", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"432.0", "achieved_ii":1, "latency":647, "max_interleaving":1, "is_fmax_bottleneck":"Yes", "is_loop_header":"Yes", "is_single_block_loop":"Yes", "loop_layer":1, "details":[{"type":"text", "text":"Variable on loop carried feedback path"}], "loop_location":{"details":[{"type":"text", "text":"Loop:  (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"58"}]}]}}, "accessMemory_0.B2":{"name":"accessMemory_0.B2", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":0, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_1.B0":{"name":"accessMemory_1.B0", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":42, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_1.B1":{"name":"accessMemory_1.B1", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"432.0", "achieved_ii":1, "latency":647, "max_interleaving":1, "is_fmax_bottleneck":"Yes", "is_loop_header":"Yes", "is_single_block_loop":"Yes", "loop_layer":1, "details":[{"type":"text", "text":"Variable on loop carried feedback path"}], "loop_location":{"details":[{"type":"text", "text":"Loop:  (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"166"}]}]}}, "accessMemory_1.B2":{"name":"accessMemory_1.B2", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":0, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_2.B0":{"name":"accessMemory_2.B0", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":42, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_2.B1":{"name":"accessMemory_2.B1", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"432.0", "achieved_ii":1, "latency":647, "max_interleaving":1, "is_fmax_bottleneck":"Yes", "is_loop_header":"Yes", "is_single_block_loop":"Yes", "loop_layer":1, "details":[{"type":"text", "text":"Variable on loop carried feedback path"}], "loop_location":{"details":[{"type":"text", "text":"Loop:  (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"274"}]}]}}, "accessMemory_2.B2":{"name":"accessMemory_2.B2", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":0, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_3.B0":{"name":"accessMemory_3.B0", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":42, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}, "accessMemory_3.B1":{"name":"accessMemory_3.B1", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"432.0", "achieved_ii":1, "latency":647, "max_interleaving":1, "is_fmax_bottleneck":"Yes", "is_loop_header":"Yes", "is_single_block_loop":"Yes", "loop_layer":1, "details":[{"type":"text", "text":"Variable on loop carried feedback path"}], "loop_location":{"details":[{"type":"text", "text":"Loop:  (%L)", "links":[{"filename":"replicated_random_access_kernels_single.cl", "line":"382"}]}]}}, "accessMemory_3.B2":{"name":"accessMemory_3.B2", "target_fmax":"Not specified", "target_ii":"Not specified", "achieved_fmax":"480.0", "achieved_ii":1, "latency":0, "max_interleaving":1, "is_fmax_bottleneck":"No", "is_loop_header":"No", "is_single_block_loop":"No", "loop_layer":0, "loop_location":{}}}, "functions":{"accessMemory_0":{"debug":[{"filename":"replicated_random_access_kernels_single.cl", "line":42}], "loop_hierachy":{"accessMemory_0__no_loop":["accessMemory_0.B0", "accessMemory_0.B2"], "accessMemory_0.B1":["accessMemory_0.B1"]}}, "accessMemory_1":{"debug":[{"filename":"replicated_random_access_kernels_single.cl", "line":150}], "loop_hierachy":{"accessMemory_1__no_loop":["accessMemory_1.B0", "accessMemory_1.B2"], "accessMemory_1.B1":["accessMemory_1.B1"]}}, "accessMemory_2":{"debug":[{"filename":"replicated_random_access_kernels_single.cl", "line":258}], "loop_hierachy":{"accessMemory_2__no_loop":["accessMemory_2.B0", "accessMemory_2.B2"], "accessMemory_2.B1":["accessMemory_2.B1"]}}, "accessMemory_3":{"debug":[{"filename":"replicated_random_access_kernels_single.cl", "line":366}], "loop_hierachy":{"accessMemory_3__no_loop":["accessMemory_3.B0", "accessMemory_3.B2"], "accessMemory_3.B1":["accessMemory_3.B1"]}}}}';
  var summaryJSON='{"performanceSummary":{"name":"Kernel Summary", "columns":["Kernel Name", "Kernel Type", "Autorun", "Workgroup Size", "# Compute Units", "Hyper-Optimized Handshaking"], "children":[{"name":"accessMemory_0", "data":["Single work-item", "No", [1, 1, 1], 1, "On"], "details":[{"type":"text", "text":"Kernel type: Single work-item"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Required workgroup size: (1, 1, 1)"}, {"type":"text", "text":"Maximum workgroup size: 1"}], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":42}]]}, {"name":"accessMemory_1", "data":["Single work-item", "No", [1, 1, 1], 1, "On"], "details":[{"type":"text", "text":"Kernel type: Single work-item"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Required workgroup size: (1, 1, 1)"}, {"type":"text", "text":"Maximum workgroup size: 1"}], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":150}]]}, {"name":"accessMemory_2", "data":["Single work-item", "No", [1, 1, 1], 1, "On"], "details":[{"type":"text", "text":"Kernel type: Single work-item"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Required workgroup size: (1, 1, 1)"}, {"type":"text", "text":"Maximum workgroup size: 1"}], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":258}]]}, {"name":"accessMemory_3", "data":["Single work-item", "No", [1, 1, 1], 1, "On"], "details":[{"type":"text", "text":"Kernel type: Single work-item"}, {"type":"text", "text":"Max global work dimension: 0"}, {"type":"text", "text":"Required workgroup size: (1, 1, 1)"}, {"type":"text", "text":"Maximum workgroup size: 1"}], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":366}]]}]}, "estimatedResources":{"name":"Estimated Resource Usage", "columns":["Kernel Name", "ALUTs ", "FFs  ", "RAMs ", "DSPs ", "MLABs"], "children":[{"name":"accessMemory_0", "data":[2264, 6282, 46, 4, 8], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":42}]]}, {"name":"accessMemory_1", "data":[2264, 6282, 46, 4, 8], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":150}]]}, {"name":"accessMemory_2", "data":[2264, 6282, 46, 4, 8], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":258}]]}, {"name":"accessMemory_3", "data":[2264, 6282, 46, 4, 8], "debug":[[{"filename":"replicated_random_access_kernels_single.cl", "line":366}]]}, {"name":"Kernel Subtotal", "classes":["summary-highlight", "nohover"], "data":[9056, 25128, 184, 16, 32]}, {"name":"Global Interconnect", "classes":["summary-highlight", "nohover"], "data":[19004, 21208, 0, 0, 0]}, {"name":"Board Interface", "classes":["summary-highlight", "nohover"], "data":[455240, 910480, 2627, 1047, 0]}, {"name":"System description ROM", "classes":["summary-highlight", "nohover"], "data":[2, 71, 2, 0, 0]}, {"name":"Total", "classes":["summary-highlight", "nohover"], "data":[483302, 956887, 2813, 1063, 32], "data_percent":[25.8971, 25.6368, 23.9997, 18.4549]}, {"name":"Available", "classes":["summary-highlight", "nohover"], "data":[1866240, 3732480, 11721, 5760, 0]}]}, "compileWarnings":{"name":"Compile Warnings", "children":[]}}';
  var infoJSON='{"name":"Info","rows":[{"classes":["info-table"],"data":["random_access_kernels_single"],"name":"Project Name"},{"data":["Stratix 10, 1SG280LU3F50E1VGS1, bitt_s10_pcie:p520_hpc_sg280l"],"name":"Target Family, Device, Board"},{"data":["19.4.0 Build 64"],"name":"AOC Version"},{"data":["19.2.0 Build 57 Pro"],"name":"Quartus Version"},{"data":["aoc /scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/device/replicated_random_access_kernels_single.cl -I/scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/device/../common -fpc -fp-relaxed -no-interleaving=default -board=p520_hpc_sg280l -o /scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/bin/random_access_kernels_single.aocx"],"name":"Command"},{"data":["Fri Apr 17 09:36:15 2020"],"name":"Reports Generated At"}]}';
  var warningsJSON='{"rows":[]}';
  var quartusJSON='{"quartusFitClockSummary":{  "name":"Quartus Fit Clock Summary"  ,"columns":["", "Kernel fmax"]  ,"children":[  {    "name":"Frequency (MHz)","data":[329.17]  }]},"quartusFitResourceUsageSummary":{  "name":"Quartus Fit Resource Utilization Summary"  ,"columns":["", "ALMs", "FFs", "RAMs", "DSPs", "MLABs"]  ,"children":[  {"name":"Full design (all kernels)","data":[10742.8 ,36869 ,168 ,16 ,32  ]  },  {"name":"accessMemory_0","data":[2689.1 ,8746 ,42 ,4 ,8]  },  {"name":"accessMemory_1","data":[2674.9 ,9352 ,42 ,4 ,8]  },  {"name":"accessMemory_2","data":[2689.6 ,9388 ,42 ,4 ,8]  },  {"name":"accessMemory_3","data":[2689.2 ,9383 ,42 ,4 ,8]  }]}}';
  var fileJSON=[{"path":"/scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/common/parameters.h", "name":"parameters.h", "has_active_debug_locs":false, "absName":"/scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/common/parameters.h", "content":"#ifndef SRC_COMMON_PARAMETERS_H_\012#define SRC_COMMON_PARAMETERS_H_\012\012/**\012 * Host specific parameters\012 */\012#define VERSION \"2.0.2\"\012#define DEFAULT_ARRAY_LENGTH 1073741824\012#define DEFAULT_PLATFORM -1\012#define DEFAULT_DEVICE -1\012#define HOST_DATA_TYPE cl_ulong\012#define HOST_DATA_TYPE_SIGNED cl_long\012#define NUM_KERNEL_REPLICATIONS 4\012#define DEFAULT_REPETITIONS 10\012\012/**\012 * Device specific parameters\012 */\012#define DEVICE_DATA_TYPE long\012#define DEVICE_DATA_TYPE_UNSIGNED ulong\012#define UNROLL_COUNT 1\012#define BUFFER_SIZE 1\012/* #undef SINGLE_KERNEL */\012// Combine address calculation and data read loop to a single loop\012#define COMBINE_LOOPS\012\012/**\012Output separator\012*/\012#define HLINE \"-------------------------------------------------------------\\n\"\012\012\012\012#endif // SRC_COMMON_PARAMETERS_H_\012"}, {"path":"/scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/device/replicated_random_access_kernels_single.cl", "name":"replicated_random_access_kernels_single.cl", "has_active_debug_locs":false, "absName":"/scratch/pc2-mitarbeiter/mariusme/synth2/frandom/cmake-build3/src/device/replicated_random_access_kernels_single.cl", "content":"/*\012Copyright (c) 2019 Marius Meyer\012\012Permission is hereby granted, free of charge, to any person obtaining a copy of\012this software and associated documentation files (the \"Software\"), to deal in\012the Software without restriction, including without limitation the rights to\012use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies\012of the Software, and to permit persons to whom the Software is furnished to do\012so, subject to the following conditions:\012\012The above copyright notice and this permission notice shall be included in all\012copies or substantial portions of the Software.\012\012THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\012IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\012FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE\012AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\012LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\012OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE\012SOFTWARE.\012*/\012\012#include \"parameters.h\"\012\012/*\012Constant used to update the pseudo random number\012*/\012#define POLY 7\012\012/*\012Kernel, that will update the given data array accoring to a predefined pseudo-\012random access scheme. The overall data array might be equally split between\012multiple kernels. In that case, the index of the current split can be given\012to the kernel.\012\012@param data The data array that will be updated\012@param m The size of the data array\012@param data_chunk The chunk index that has to be updated by the kernel\012*/\012__attribute__((max_global_work_dim(0)))\012__kernel\012void accessMemory_0(__global DEVICE_DATA_TYPE_UNSIGNED  volatile * restrict data,\012                        const DEVICE_DATA_TYPE_UNSIGNED m,\012                        const DEVICE_DATA_TYPE_UNSIGNED data_chunk,\012                        const uint kernel_number) {\012    // Initiate the pseudo random number\012    DEVICE_DATA_TYPE_UNSIGNED ran = 1;\012\012    // calculate the start of the address range this kernel is responsible for\012    #ifndef SINGLE_KERNEL\012    DEVICE_DATA_TYPE_UNSIGNED const address_start = kernel_number * data_chunk;\012    #endif\012\012    DEVICE_DATA_TYPE_UNSIGNED const mupdate = 4 * m;\012\012    // do random accesses\012#pragma ivdep\012    for (DEVICE_DATA_TYPE_UNSIGNED i=0; i< mupdate / BUFFER_SIZE; i++) {\012\012        DEVICE_DATA_TYPE_UNSIGNED local_address[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED loaded_data[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED update_val[BUFFER_SIZE];\012\012        // calculate next addresses\012        __attribute__((opencl_unroll_hint(1)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            DEVICE_DATA_TYPE v = 0;\012            if (((DEVICE_DATA_TYPE) ran) < 0) {\012                v = POLY;\012            }\012            ran = (ran << 1) ^ v;\012            update_val[ld] = ran;\012            DEVICE_DATA_TYPE_UNSIGNED address = (ran >> 3) & (m - 1);\012            #ifndef SINGLE_KERNEL\012            local_address[ld] = address - address_start;\012            #else\012            local_address[ld] = address;\012            #endif\012#ifndef COMBINE_LOOPS\012        }\012\012        // load the data of the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012#endif\012            #ifdef SINGLE_KERNEL\012            loaded_data[ld] = data[local_address[ld]];\012            #else\012            if (local_address[ld] < data_chunk) {\012                loaded_data[ld] = data[local_address[ld]];\012            }\012            #endif\012        }\012\012        // store back the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            #ifdef SINGLE_KERNEL\012            data[local_address[ld]] = loaded_data[ld] ^update_val[ld];\012            #else\012            if (local_address[ld] < data_chunk) {\012                data[local_address[ld]] = loaded_data[ld] ^ update_val[ld];\012            }\012            #endif\012        }\012    }\012}\012\012/*\012Copyright (c) 2019 Marius Meyer\012\012Permission is hereby granted, free of charge, to any person obtaining a copy of\012this software and associated documentation files (the \"Software\"), to deal in\012the Software without restriction, including without limitation the rights to\012use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies\012of the Software, and to permit persons to whom the Software is furnished to do\012so, subject to the following conditions:\012\012The above copyright notice and this permission notice shall be included in all\012copies or substantial portions of the Software.\012\012THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\012IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\012FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE\012AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\012LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\012OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE\012SOFTWARE.\012*/\012\012#include \"parameters.h\"\012\012/*\012Constant used to update the pseudo random number\012*/\012#define POLY 7\012\012/*\012Kernel, that will update the given data array accoring to a predefined pseudo-\012random access scheme. The overall data array might be equally split between\012multiple kernels. In that case, the index of the current split can be given\012to the kernel.\012\012@param data The data array that will be updated\012@param m The size of the data array\012@param data_chunk The chunk index that has to be updated by the kernel\012*/\012__attribute__((max_global_work_dim(0)))\012__kernel\012void accessMemory_1(__global DEVICE_DATA_TYPE_UNSIGNED  volatile * restrict data,\012                        const DEVICE_DATA_TYPE_UNSIGNED m,\012                        const DEVICE_DATA_TYPE_UNSIGNED data_chunk,\012                        const uint kernel_number) {\012    // Initiate the pseudo random number\012    DEVICE_DATA_TYPE_UNSIGNED ran = 1;\012\012    // calculate the start of the address range this kernel is responsible for\012    #ifndef SINGLE_KERNEL\012    DEVICE_DATA_TYPE_UNSIGNED const address_start = kernel_number * data_chunk;\012    #endif\012\012    DEVICE_DATA_TYPE_UNSIGNED const mupdate = 4 * m;\012\012    // do random accesses\012#pragma ivdep\012    for (DEVICE_DATA_TYPE_UNSIGNED i=0; i< mupdate / BUFFER_SIZE; i++) {\012\012        DEVICE_DATA_TYPE_UNSIGNED local_address[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED loaded_data[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED update_val[BUFFER_SIZE];\012\012        // calculate next addresses\012        __attribute__((opencl_unroll_hint(1)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            DEVICE_DATA_TYPE v = 0;\012            if (((DEVICE_DATA_TYPE) ran) < 0) {\012                v = POLY;\012            }\012            ran = (ran << 1) ^ v;\012            update_val[ld] = ran;\012            DEVICE_DATA_TYPE_UNSIGNED address = (ran >> 3) & (m - 1);\012            #ifndef SINGLE_KERNEL\012            local_address[ld] = address - address_start;\012            #else\012            local_address[ld] = address;\012            #endif\012#ifndef COMBINE_LOOPS\012        }\012\012        // load the data of the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012#endif\012            #ifdef SINGLE_KERNEL\012            loaded_data[ld] = data[local_address[ld]];\012            #else\012            if (local_address[ld] < data_chunk) {\012                loaded_data[ld] = data[local_address[ld]];\012            }\012            #endif\012        }\012\012        // store back the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            #ifdef SINGLE_KERNEL\012            data[local_address[ld]] = loaded_data[ld] ^update_val[ld];\012            #else\012            if (local_address[ld] < data_chunk) {\012                data[local_address[ld]] = loaded_data[ld] ^ update_val[ld];\012            }\012            #endif\012        }\012    }\012}\012\012/*\012Copyright (c) 2019 Marius Meyer\012\012Permission is hereby granted, free of charge, to any person obtaining a copy of\012this software and associated documentation files (the \"Software\"), to deal in\012the Software without restriction, including without limitation the rights to\012use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies\012of the Software, and to permit persons to whom the Software is furnished to do\012so, subject to the following conditions:\012\012The above copyright notice and this permission notice shall be included in all\012copies or substantial portions of the Software.\012\012THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\012IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\012FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE\012AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\012LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\012OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE\012SOFTWARE.\012*/\012\012#include \"parameters.h\"\012\012/*\012Constant used to update the pseudo random number\012*/\012#define POLY 7\012\012/*\012Kernel, that will update the given data array accoring to a predefined pseudo-\012random access scheme. The overall data array might be equally split between\012multiple kernels. In that case, the index of the current split can be given\012to the kernel.\012\012@param data The data array that will be updated\012@param m The size of the data array\012@param data_chunk The chunk index that has to be updated by the kernel\012*/\012__attribute__((max_global_work_dim(0)))\012__kernel\012void accessMemory_2(__global DEVICE_DATA_TYPE_UNSIGNED  volatile * restrict data,\012                        const DEVICE_DATA_TYPE_UNSIGNED m,\012                        const DEVICE_DATA_TYPE_UNSIGNED data_chunk,\012                        const uint kernel_number) {\012    // Initiate the pseudo random number\012    DEVICE_DATA_TYPE_UNSIGNED ran = 1;\012\012    // calculate the start of the address range this kernel is responsible for\012    #ifndef SINGLE_KERNEL\012    DEVICE_DATA_TYPE_UNSIGNED const address_start = kernel_number * data_chunk;\012    #endif\012\012    DEVICE_DATA_TYPE_UNSIGNED const mupdate = 4 * m;\012\012    // do random accesses\012#pragma ivdep\012    for (DEVICE_DATA_TYPE_UNSIGNED i=0; i< mupdate / BUFFER_SIZE; i++) {\012\012        DEVICE_DATA_TYPE_UNSIGNED local_address[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED loaded_data[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED update_val[BUFFER_SIZE];\012\012        // calculate next addresses\012        __attribute__((opencl_unroll_hint(1)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            DEVICE_DATA_TYPE v = 0;\012            if (((DEVICE_DATA_TYPE) ran) < 0) {\012                v = POLY;\012            }\012            ran = (ran << 1) ^ v;\012            update_val[ld] = ran;\012            DEVICE_DATA_TYPE_UNSIGNED address = (ran >> 3) & (m - 1);\012            #ifndef SINGLE_KERNEL\012            local_address[ld] = address - address_start;\012            #else\012            local_address[ld] = address;\012            #endif\012#ifndef COMBINE_LOOPS\012        }\012\012        // load the data of the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012#endif\012            #ifdef SINGLE_KERNEL\012            loaded_data[ld] = data[local_address[ld]];\012            #else\012            if (local_address[ld] < data_chunk) {\012                loaded_data[ld] = data[local_address[ld]];\012            }\012            #endif\012        }\012\012        // store back the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            #ifdef SINGLE_KERNEL\012            data[local_address[ld]] = loaded_data[ld] ^update_val[ld];\012            #else\012            if (local_address[ld] < data_chunk) {\012                data[local_address[ld]] = loaded_data[ld] ^ update_val[ld];\012            }\012            #endif\012        }\012    }\012}\012\012/*\012Copyright (c) 2019 Marius Meyer\012\012Permission is hereby granted, free of charge, to any person obtaining a copy of\012this software and associated documentation files (the \"Software\"), to deal in\012the Software without restriction, including without limitation the rights to\012use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies\012of the Software, and to permit persons to whom the Software is furnished to do\012so, subject to the following conditions:\012\012The above copyright notice and this permission notice shall be included in all\012copies or substantial portions of the Software.\012\012THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\012IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\012FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE\012AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\012LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\012OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE\012SOFTWARE.\012*/\012\012#include \"parameters.h\"\012\012/*\012Constant used to update the pseudo random number\012*/\012#define POLY 7\012\012/*\012Kernel, that will update the given data array accoring to a predefined pseudo-\012random access scheme. The overall data array might be equally split between\012multiple kernels. In that case, the index of the current split can be given\012to the kernel.\012\012@param data The data array that will be updated\012@param m The size of the data array\012@param data_chunk The chunk index that has to be updated by the kernel\012*/\012__attribute__((max_global_work_dim(0)))\012__kernel\012void accessMemory_3(__global DEVICE_DATA_TYPE_UNSIGNED  volatile * restrict data,\012                        const DEVICE_DATA_TYPE_UNSIGNED m,\012                        const DEVICE_DATA_TYPE_UNSIGNED data_chunk,\012                        const uint kernel_number) {\012    // Initiate the pseudo random number\012    DEVICE_DATA_TYPE_UNSIGNED ran = 1;\012\012    // calculate the start of the address range this kernel is responsible for\012    #ifndef SINGLE_KERNEL\012    DEVICE_DATA_TYPE_UNSIGNED const address_start = kernel_number * data_chunk;\012    #endif\012\012    DEVICE_DATA_TYPE_UNSIGNED const mupdate = 4 * m;\012\012    // do random accesses\012#pragma ivdep\012    for (DEVICE_DATA_TYPE_UNSIGNED i=0; i< mupdate / BUFFER_SIZE; i++) {\012\012        DEVICE_DATA_TYPE_UNSIGNED local_address[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED loaded_data[BUFFER_SIZE];\012        DEVICE_DATA_TYPE_UNSIGNED update_val[BUFFER_SIZE];\012\012        // calculate next addresses\012        __attribute__((opencl_unroll_hint(1)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            DEVICE_DATA_TYPE v = 0;\012            if (((DEVICE_DATA_TYPE) ran) < 0) {\012                v = POLY;\012            }\012            ran = (ran << 1) ^ v;\012            update_val[ld] = ran;\012            DEVICE_DATA_TYPE_UNSIGNED address = (ran >> 3) & (m - 1);\012            #ifndef SINGLE_KERNEL\012            local_address[ld] = address - address_start;\012            #else\012            local_address[ld] = address;\012            #endif\012#ifndef COMBINE_LOOPS\012        }\012\012        // load the data of the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012#endif\012            #ifdef SINGLE_KERNEL\012            loaded_data[ld] = data[local_address[ld]];\012            #else\012            if (local_address[ld] < data_chunk) {\012                loaded_data[ld] = data[local_address[ld]];\012            }\012            #endif\012        }\012\012        // store back the calculated addresses from global memory\012        __attribute__((opencl_unroll_hint(UNROLL_COUNT)))\012        for (int ld=0; ld< BUFFER_SIZE; ld++) {\012            #ifdef SINGLE_KERNEL\012            data[local_address[ld]] = loaded_data[ld] ^update_val[ld];\012            #else\012            if (local_address[ld] < data_chunk) {\012                data[local_address[ld]] = loaded_data[ld] ^ update_val[ld];\012            }\012            #endif\012        }\012    }\012}\012\012"}];
  