Forum Discussion
New ContributorStable argument doesn't work in simulation
Thanks for sharing the report, @DorianL .
It looks like the loop at line 137 was pipelined with II=1, but it was constrained to serial execution.
This means that this outer loop is effectively un-pipelined. This doesn't explain the gaps you are seeing in the simulation waveform though.
I also see that you are getting a memory system with lots of arbitration:After some experimenting, I discovered that the warning about the variable 'fenetre' is a bit of a red herring here. I would expect for a line buffer like you are describing to have a memory system with multiple banks, and each bank having a dedicated load/store unit (LSU). From the image above, we can see that the memory system is not efficiently selecting banks. I tried using the bank_bits attribute to constrain this, but it appears the compiler is ignoring this attribute now.
I was able to get the compiler to partition your 2d array by swapping the dimensions (transposing) so that the dimension to be split into banks (i.e. accessed simultaneously by different unrolled loop iterations) was in the least significant place. This appears to result in the desired memory system (don't forget to swap the accesses too!!)OLD:
// Ligne a retard [[intel::fpga_memory("BLOCK_RAM")]] // memory unsigned int line_buffer[8][NB_COLONNE_MAX];NEW:
// Ligne a retard [[intel::fpga_memory("BLOCK_RAM")]] // memory unsigned int line_buffer[NB_COLONNE_MAX][8];* Note that I changed the dimension from 5 to 8: the compiler complains if you try to create a memory system with a non-power-of-2 number of banks. Changing to 8 is ok because the compiler sees that the extra 3 banks aren't used and it optimizes them away.
The new memory system looks a lot better now:
The sim looks a lot better too:
I think i know how to solve these 2-cycle dips but I'm still waiting for the test to finish.
I suspect it's a side-effect of using a loop nest instead of using a while(1) loop to iterate across image pixels.
ModeratorHi @DorianL,
Greetings, just to further understand the issues you mention that you have a inner loop issues.
Would you be able to share the code example that you have for the kernels and also what are the error/issues that you are seeing? And is it correct to assume that the error happens on the compilation?
If you can provide the compilation command it would also be very helpful.
That would better help us to understand the situation.
Thanks,
Regards
BB
New Contributortemplate <typename flux_in, typename flux_tempo> struct travail_sur_voisinage { sycl::ext::oneapi::experimental::annotated_arg< int , decltype(sycl::ext::oneapi::experimental::properties{ stable})> taille_h; sycl::ext::oneapi::experimental::annotated_arg< int , decltype(sycl::ext::oneapi::experimental::properties{ stable})> taille_v; auto get(sycl::ext::oneapi::experimental::properties_tag) { return sycl::ext::oneapi::experimental::properties{ streaming_interface<>}; } void operator()() const { //Compteurs ligne pixel //Entree Sortie [[intel::fpga_register]] unsigned int pixel_a_traiter; [[intel::fpga_register]] unsigned int pixel_a_envoyer; [[intel::fpga_register]] unsigned int pixel_apres_traitement; //Ligne a retard [[intel::fpga_memory("BLOCK_RAM")]] unsigned int line_buffer[5][NB_COLONNE_MAX]; //Voisinnage [[intel::fpga_register]] unsigned int fenetre[5][5]; [[intel::initiation_interval(1)]] for (int num_lig = 0; num_lig < taille_v + 2; num_lig++) { [[intel::initiation_interval(1)]] //[[intel::speculated_iterations(0)]] for (int num_col = 0; (num_col < taille_h + 2); num_col++) { if (num_lig < taille_v && num_col < taille_h) { pixel_a_traiter = flux_in::read(); //Gestion ligne a retard fpga_tools::UnrolledLoop<0,4>([&](auto l) { line_buffer[l][num_col] = line_buffer[l + 1][num_col]; }); line_buffer[4][num_col] = pixel_a_traiter; //Fin gestion ligne a retard //Fenetre video glissante fpga_tools::UnrolledLoop<0,5>([&](auto li) { // #pragma unroll fpga_tools::UnrolledLoop<0,4>([&](auto co) { fenetre[li][co] = fenetre[li][co + 1]; }); fenetre[li][4] = line_buffer[li][num_col]; }); //Fin Fenetre video glissante } pixel_apres_traitement = traitement_5x5(fenetre); if ((num_lig >= 2) && (num_col >= 2)) { pixel_a_envoyer = 0; if (((num_lig >= 4) && (num_lig < taille_v) && (num_col >= 4) && (num_col < taille_h))) { pixel_a_envoyer = pixel_apres_traitement; } flux_tempo::write(pixel_a_envoyer); } } } } };
the compilation command is :
tp3_video.fpga_sim: kernel_sim.o icpx -fsycl -fintelfpga -Xsclock=400MHz -Xsoptimize=latency -Xssimulation -Xsghdl=0 -Xstarget=Agilex7 -Xsv $^ -o $@ kernel_sim.o: src/tp3_video.cpp icpx -fsycl -O3 -g -std=c++17 -Wall -I include -v -fintelfpga -Xsprofile -Xssimulation -DFPGA_SIMULATOR -o $@ -c src/tp3_video.cpp
Have a great day,
DorianL
