9 Replies Latest reply on Sep 15, 2010 10:29 AM by himanshu.gautam

    Can non-kernel functions access global memory?

    edwen

      My .cl is somehow like this:

      void evaluation(float *L)

      {...}

      __kernel void memdata(__global float *d_v, __global float LC, const int num)

      { const int tid = get_global_id(0);

      __global float *L_loc = LC + num*tid;

      evaluation(L_loc);

       }

       

      This always brings me a programbuild error. The reason I have to add "__global" in front of L_loc is because LC is in global memory, otherwise L_loc can't point to LC at all. However, "evaluation" is a non-kernel function, and it seems it can't reach global memory. Is that right?

        • Can non-kernel functions access global memory?
          n0thing

          This compiles fine after changing the argument pointer of the function to point to global memory.

          void evaluation(__global float *L)

          {}

          __kernel void memdata(__global float *d_v, __global float *LC, const int num)

          { const int tid = get_global_id(0);

          __global float *L_loc = LC + num*tid;

          evaluation(L_loc);

           }

            • Can non-kernel functions access global memory?
              edwen

              Yes, I tried that solution, too. It compiles fine, but fails when I try to run it.

                • Can non-kernel functions access global memory?
                  n0thing

                  Fails as in correctness failure? Can you post the complete code?

                  BTW The following code works just fine for me -

                  unsigned int func(__global unsigned int *ptr, unsigned int mul) { return ptr[get_global_id(0)] * mul; } __kernel void templateKernel(__global unsigned int * output, __global unsigned int * input, const unsigned int multiplier) { uint tid = get_global_id(0); output[tid] = func(input,multiplier); }

                • Can non-kernel functions access global memory?
                  edwen

                   

                  Originally posted by: n0thing This compiles fine after changing the argument pointer of the function to point to global memory.

                  void evaluation(__global float *L)

                  Seems like in a non-kernel function, it's not allow to add __global in front a parameter. You can try to add any irrelevant function with that kind of parameter in you own code and see. Thanks,

                   

                    • Can non-kernel functions access global memory?
                      karbous

                      I can confirm I'm using __global parameters in non-kernel functions successfully. So maybe try out changing your code to something following: 

                      void evaluation(__global float *L, const int startIndex) { //acces the array as L[startIndex] } __kernel void memdata(__global float *d_v, __global float *LC, const int num) { const int tid = get_global_id(0); evaluation(LC, num*tid); }

                        • Can non-kernel functions access global memory?
                          edwen

                          Thanks, n0thing and karbous. Yes, n0thing's code works fine. I will figure out what's wrong with mine. I couldn't post the whole thing before i get the permission from the project manager.

                          • Can non-kernel functions access global memory?
                            edwen

                             

                            Originally posted by: karbous I can confirm I'm using __global parameters in non-kernel functions successfully. So maybe try out changing your code to something following: 

                            I think I figured out what my problem is. Please look at the code:


                            float valuation(const int t, float *L, const int numFactors_, const int numLibors_, const int numTimeSteps_, const InstrumentType instrumentType_)
                            {float *npv;
                            float npv_;
                            npv_ = 0.0f;
                            npv = &npv_;

                            if (instrumentType_ == SWAP)
                            {swap_valuation(t, L, npv);}

                            else if (instrumentType_ == CAP)
                            {cap_valuation(t, L, npv);}

                            else if (instrumentType_ == SWAPTION)
                            {swaption_valuation(t, L, npv, numFactors_, numLibors_, numTimeSteps_);
                            }
                            return *npv;
                            }

                            __kernel void fun1(__global float *d_v,...)
                            {...
                            evaluation(...);
                            }

                            __kernel void fun2(__global float *d_v, __global float *L,...)
                            {...
                            evaluation(...);
                            }

                            As in above, I have two kernel functions, and they both call the "evalution" function. However, in kernel one, I dont have a global array argument named "L" passed from my host code, and the "L" array required by function "evaluation" will be generated locally. In the second kernel function, I do have an global argument named "L" passed from the host code. If I add "__global" to the parameter "L" in function "evaluation", the first kernel function wont work any more. Further more, all the functions inside "evaluation" would have problems, too, since they dont have a "__global" sign, either.

                             

                            OpenCL should really allows overriding.