Altera_Forum
Honored Contributor
Honored ContributorHow to improve memory bandwidth
Hi, I want to do empty kernels.
First block read data from memory, send this data to channel. Second block receive data from channel and write this data to memory. I execute this code on Arria10 DeVKit But information of bandwidth in profile is just ~3000 Mb/sec, but max bandwidth is 12800 Mb/sec How can I improve this result ?
# pragma OPENCL_EXTENSION cl_altera_channels : enable
# define THREAD_SIZE 1# define VEC_SIZE 128
typedef struct{
float data;
} lane_data;
channel lane_data data_ch;
__kernel
void memRead(
int data_dim1,
int data_dim2,
int data_dim3,
__global float *restrict bottom
)
{
lane_data buff;
int size = data_dim1*data_dim2*data_dim3;
for(int ll=0; ll<size/THREAD_SIZE/VEC_SIZE; ll++){
# pragma unroll
for(int i=0; i<THREAD_SIZE; i++){
# pragma unroll
for(int j=0; j<VEC_SIZE; j++){
buff.data = bottom;
}
write_channel_altera(data_ch, buff);
}
}
}
__kernel
void memWrite(
int data_dim1,
int data_dim2,
int data_dim3,
__global float *restrict result
)
{
lane_data buff;
int size = data_dim1*data_dim2*data_dim3;
for(int ll=0; ll<size/THREAD_SIZE/VEC_SIZE; ll++){
# pragma unroll
for(int i=0; i<THREAD_SIZE; i++){
buff = read_channel_altera(data_ch);
# pragma unroll
for(int j=0; j<VEC_SIZE; j++){
result = buff.data;
}
}
}
}
