11 Replies Latest reply on May 25, 2011 11:01 AM by donnervogel80

    how to initialize a constant variable of array?

    edwen
      constant memory initialization

      I am trying to modify a CUDA program into OpenCL and I met a problem. In my "market.cl" file, I need to declare and initialize some constant variables. for example, "__constant int v1 = 5;" This works fine. However, I need to declare some arrays, such as "__constant float delta[5][80]". In my CUDA program, I used the function "cudaMemcpyToSymbol" to copy all constants into constant memory, and the code in the "market.cu" is something like:

      #define NumFactors 8

      float delta_[5][80];

      float weight = sqrt(1.0f/float)NumFactors);

      for (int i = 0; i < 5; i++)

      {  for (int j = 0; j < 80; j++)

               delta_[j] = weight;

      }

      cudaMemcpyToSymbol (delta, &delta_, sizeof(delta_);

       

      How can I realize this in OpenCL? Thank you very much for your advice.

       

        • how to initialize a constant variable of array?
          genaganna

           

          Originally posted by: edwen I am trying to modify a CUDA program into OpenCL and I met a problem. In my "market.cl" file, I need to declare and initialize some constant variables. for example, "__constant int v1 = 5;" This works fine. However, I need to declare some arrays, such as "__constant float delta[5][80]". In my CUDA program, I used the function "cudaMemcpyToSymbol" to copy all constants into constant memory, and the code in the "market.cu" is something like:

          #define NumFactors 8

          float delta_[5][80];

          float weight = sqrt(1.0f/float)NumFactors);

          for (int i = 0; i < 5; i++)

          {  for (int j = 0; j < 80; j++)

                   delta_[j] = weight;

          }

          cudaMemcpyToSymbol (delta, &delta_, sizeof(delta_);

          How can I realize this in OpenCL? Thank you very much for your advice.

           

           



          Setting __constant variable array from runtime is exactly same as global memory variable.

          1. Create a buffer and assign values from delta_ array.

          2. Set kernel argument to appropriate kernel parameter.

            • how to initialize a constant variable of array?
              edwen

              Setting __constant variable from running is exactly same as global memory.

              1. Create a buffer and assign values from delta_ array.

              2. Set kernel argument to appropriate kernel parameter.

              Thanks for your prompt reply. However, in my "market.cl", besides two kernel functions, there are several other functions. I need to use the "delta" array in these functions, and then the two kernel functions call these functions.  Therefore, I can not just pass the array as argument to certain kernel parameter. The structure of my "market.cl" is somehow like:

              float func1(...)

              { ... //code using array "delta"}

              float func2(...)

              {.../code using array "delta"}

              __kernel func (...)

              {...;

              y1 = func1(...);

              y2 = func2(...)}

              there would be a run-time error if I pass array "delta" to the kernel function "func" as argument since "func1" and "func2" need this array.

               

                • how to initialize a constant variable of array?
                  genaganna

                   

                  Originally posted by: edwen


                   

                  Setting __constant variable from running is exactly same as global memory.

                   

                  1. Create a buffer and assign values from delta_ array.

                   

                  2. Set kernel argument to appropriate kernel parameter.

                   

                   

                  Thanks for your prompt reply. However, in my "market.cl", besides two kernel functions, there are several other functions. I need to use the "delta" array in these functions, and then the two kernel functions call these functions.  Therefore, I can not just pass the array as argument to certain kernel parameter. The structure of my "market.cl" is somehow like:

                   

                  float func1(...)

                   

                  { ... //code using array "delta"}

                   

                  float func2(...)

                   

                  {.../code using array "delta"}

                   

                  __kernel func (...)

                   

                  {...;

                   

                  y1 = func1(...);

                   

                  y2 = func2(...)}

                   

                  there would be a run-time error if I pass array "delta" to the kernel function "func" as argument since "func1" and "func2" need this array.

                   

                   

                  Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

                  __constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

                   

                    • how to initialize a constant variable of array?
                      edwen

                       

                      Originally posted by: genaganna
                      Originally posted by: edwen


                      Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

                      __constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

                       

                      I thought about this way. However, too reasons this doesn't work for me that well. Firstly, the array is very big, and I need to use a for loop to generate the values. Secondly, in the for loop generating the values of the array element, some math functions, such as sqrt and exp, are used. It's not applicable to initialize the array using your method. I wonder if there is an alternative way in OpenCL. Thanks.

                        • how to initialize a constant variable of array?
                          genaganna

                           

                          Originally posted by: edwen
                          Originally posted by: genaganna
                          Originally posted by: edwen


                           

                          Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

                           

                          __constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

                           

                           

                           

                           

                          I thought about this way. However, too reasons this doesn't work for me that well. Firstly, the array is very big, and I need to use a for loop to generate the values. Secondly, in the for loop generating the values of the array element, some math functions, such as sqrt and exp, are used. It's not applicable to initialize the array using your method. I wonder if there is an alternative way in OpenCL. Thanks.

                           

                          1. Initialize values of delta_ in host and create CLBuffer and assign this values to CLBuffer

                          2. All kernels should have a __constant parameter so that you can pass your value to all kernels.

                           

                          what is size of your __constant buffer size.

                            • how to initialize a constant variable of array?
                              edwen

                               

                              Originally posted by: genaganna
                              Originally posted by: edwen
                              Originally posted by: genaganna
                              Originally posted by: edwen


                              1. Initialize values of delta_ in host and create CLBuffer and assign this values to CLBuffer

                              2. All kernels should have a __constant parameter so that you can pass your value to all kernels.

                               

                              what is size of your __constant buffer size.

                              The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

                                • how to initialize a constant variable of array?
                                  genaganna

                                   

                                  Originally posted by: edwen

                                   

                                  The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

                                   

                                  I said all kernels should have this as parameter.

                                    • how to initialize a constant variable of array?
                                      edwen

                                       

                                      Originally posted by: genaganna
                                      Originally posted by: edwen

                                       

                                      The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

                                       

                                      I said all kernels should have this as parameter.

                                      ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right? Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

                                        • how to initialize a constant variable of array?
                                          genaganna

                                           

                                          Originally posted by: edwen

                                           

                                          ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right?

                                          Yes you are right.

                                          Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

                                           

                                          Yes, it is not possible to pass any arguments to non-kernel functions.

                                            • how to initialize a constant variable of array?
                                              edwen

                                               

                                              Originally posted by: genaganna
                                              Originally posted by: edwen

                                               

                                              ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right?

                                              Yes you are right.

                                              Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

                                               

                                              Yes, it is not possible to pass any arguments to non-kernel functions.

                                              This method is so inefficient. I have more than 10 constant variables in my .cl file, which are used by those non-kernal functions. Now I have to add corresonding parameters to all these functions. And accordingly, in my kernel functions, add extra parameters while calling these non-kernel functions. I think the CUDA method is much better, namely we only declare the constant variables in .cl file, but calculate the values of these variables in the host, and then copy them into the constant memory through sth similar to cudaMemcpyToSymbol(). OpenCL should improve the way it declares and initializes variables in constant memories.

                                                • how to initialize a constant variable of array?
                                                  donnervogel80

                                                  There is one other possibility, in case you are using run-time compilation from source in your opencl code. You can generate the code for initializing the constant varible in the main program, add it to the kernel code, and then compile it. Once you have all your data for the constants, you can for sure write a function which generates this kind of code as a string:

                                                  __constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

                                                  Then prefix your opencl code with this string and compile.

                                                  OpenCL is run-time compiled - use that for your best advantage! Think more about that... you can replace parameters with compile-time defines (allowing the compiler to unroll loops), etc.