- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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 = read_channel_intel(c0); //read from the channel}
for(int i = 0; i < n; i ++){
//do some computation of the page_rank 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?
Link Copied
3 Replies
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
Are you unrolling your loops? Post your complete kernel since the snippet you have posted should compile correctly.
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
--- Quote Start --- Are you unrolling your loops? Post your complete kernel since the snippet you have posted should compile correctly. --- Quote End --- 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 = 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 = 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); } }
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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.
Reply
Topic Options
- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page