Stable argument doesn't work in simulation

Hi @aikeu,

Thank you for your reply, I already checked the documentation and I think the issue is that my inner loop is dynamic because it is an input of my IP (the number of column of the picture to process). I tried to solve the issue by bounding the loop with a random number like it is made in the example of the documentation, but my kernel doesn't accept the rand() function because it's an external SYCL function. Do you have an idea of how I could make my kernel work ?

Thank you !

DorianL

Hi @BoonBengT_Altera,

template <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

Hi @aikeu,

Thank you for your help, I'll check that now !

DorianL

Hi @aikeu ,

I was waiting for an answer following my last post with my code and compilation commands as asked by @BoonBengT_Altera in a previous post. Could we wait an answer from him before closing the thread ? Thank you !

DorianL

Hi @whitepau_altera,

Thank you for your answer, I tried to bound the loop with a constant number and it doesn't change anything on the simulation, the read_rdy signal of the pipe is sometimes low whereas it is always high when i use a constant size of image instead of kernel arguments.

The issue is still the same the inner loop is very slow when I compare between kernel argument and constant value.

I tried to coalesce the loop nest but I have worse results in simulation than when I don't use coalesce loops.

I already tried to use the convolution_2d github example and I can't use a picture with higher dimension than the original example (test_0.bmp,test_1.bmp...). Is this possible to use other pictures ?

Thank you,

DorianL

Hi @whitepau_altera,

Thank you for your work and answer, everything is working good now ! I never thought it could be a memory issue on that, I will check the kernel memory viewer more often now. The simulation is perfect now and the rdy signal is as it should be. Have a great day !

DorianL

Hi @whitepau_altera,

Of course you can use it as an example for kernel memory optimization. Thank you !

DorianL