Results 1 to 4 of 4

Thread: writing and reading the same pipe in one kernel

  1. #1
    Join Date
    May 2018
    Posts
    2
    Rep Power
    1

    Default writing and reading the same pipe in one kernel

    Hi
    I want to write and read the same pipe in one kernel like this:

    channel float c0 __attribute__((depth(2)));
    __kernel void testkernel() {

    for(int i = 0; i < n; i ++){
    page_ranks[i] = read_channel_intel(c0); //read from the channel
    }

    for(int i = 0; i < n; i ++){
    //do some computation of the page_rank[i] and generage new_rank
    write_channel_intel(c0, new_rank); //write new_rank to this pipe

    }

    }
    and this kernel is being run in multiple iterations, each time it will read from last iteration's result and compute new rank and use the new rank as the new input for the next iteration.
    When I compile this, I get the error of "Compiler Error: Multiple writes to channel c0", this is the only line I got in the log.
    Any suggestions on how to deal with this case using pipe or how to resolve this bug?

  2. #2
    Join Date
    Jan 2017
    Posts
    692
    Rep Power
    1

    Default Re: writing and reading the same pipe in one kernel

    Are you unrolling your loops? Post your complete kernel since the snippet you have posted should compile correctly.

  3. #3
    Join Date
    May 2018
    Posts
    2
    Rep Power
    1

    Default Re: writing and reading the same pipe in one kernel

    Quote Originally Posted by HRZ View Post
    Are you unrolling your loops? Post your complete kernel since the snippet you have posted should compile correctly.
    Below is my kernel code:

    #define M 100
    channel float c0 __attribute__((depth(2)));


    __kernel void mapreduce_page_rank(__global float* restrict page_ranks,
    int n,
    const __global int* restrict pages,
    const __global unsigned int* restrict noutlinks){


    int i, j, t;
    float new_rank;
    float d_factor = 0.85;
    float cur;
    float maps;
    float outbound_rank;
    float new_rank_copies[M];


    #pragma unroll
    for (i = 0; i < M; i ++){
    new_rank_copies[i] = 0;
    }


    for(i=0; i<n; ++i){
    new_rank = 0.0;
    for(j=0; j<n; ++j){
    outbound_rank = page_ranks[j] / (double)noutlinks[j];
    maps = pages[i*n+j] * outbound_rank;
    cur = new_rank_copies[M-1] + maps;
    #pragma unroll
    for(int c = M-1; c > 0; c--){
    new_rank_copies[c] = new_rank_copies[c - 1];
    }
    new_rank_copies[0] = cur;
    }
    #pragma unroll
    for(int j = 0; j < M; j ++){
    new_rank += new_rank_copies[j];
    new_rank_copies[j] = 0;
    }


    new_rank = ((1-d_factor)/n)+(d_factor*new_rank);
    write_channel_altera(c0, new_rank);
    }
    }




    __kernel void produce_page_rank(__global float* restrict page_ranks,
    int n,
    const __global int* restrict pages,
    const __global unsigned int* restrict noutlinks){


    int i, j, t;
    float new_rank;
    float d_factor = 0.85;
    float cur;
    float maps;
    float outbound_rank;
    float new_rank_copies[M];


    #pragma unroll
    for (i = 0; i < M; i ++){
    new_rank_copies[i] = 0;
    }


    for(int i = 0; i < n; i ++){
    page_ranks[i] = read_channel_altera(c0);
    }


    for(i=0; i<n; ++i){
    new_rank = 0.0;
    for(j=0; j<n; ++j){
    outbound_rank = page_ranks[j] / (double)noutlinks[j];
    maps = pages[i*n+j] * outbound_rank;
    cur = new_rank_copies[M-1] + maps;
    #pragma unroll
    for(int c = M-1; c > 0; c--){
    new_rank_copies[c] = new_rank_copies[c - 1];
    }
    new_rank_copies[0] = cur;
    }
    #pragma unroll
    for(int j = 0; j < M; j ++){
    new_rank += new_rank_copies[j];
    new_rank_copies[j] = 0;
    }


    new_rank = ((1-d_factor)/n)+(d_factor*new_rank);
    write_channel_altera(c0, new_rank);
    }
    }

  4. #4
    Join Date
    Jan 2017
    Posts
    692
    Rep Power
    1

    Default Re: writing and reading the same pipe in one kernel

    Well, you are clearly writing to the same channel twice, once in your first kernel and the second time in your second kernel. This is not allowed since each channel works as a FIFO with one read point and one write point. However, I think the newer versions of the compiler (v17 or v17.1) allow multiple channel call sites but I haven't tried it myself.

Similar Threads

  1. Replies: 7
    Last Post: October 1st, 2015, 11:42 PM
  2. reading and writing to ddr2
    By jbouie in forum General Altera Discussion
    Replies: 2
    Last Post: September 25th, 2012, 11:54 PM
  3. These is something wrong about DDR reading and writing
    By wangheng0987654321 in forum General Discussion Forum
    Replies: 0
    Last Post: December 19th, 2011, 12:59 AM
  4. DE1 Reading / Writing to SRAM
    By jtf323 in forum Quartus II and EDA Tools Discussion
    Replies: 4
    Last Post: January 13th, 2011, 03:06 AM
  5. Writing and Reading Data To & From RAM
    By digitallogic in forum FPGA, Hardcopy, and CPLD Discussion
    Replies: 2
    Last Post: December 8th, 2010, 01:59 AM

Tags for this Thread

Bookmarks

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •