4 Replies Latest reply on Oct 5, 2009 2:06 PM by ochafik

    2 dimensional fixed size arrays supported?

    oscarbarenys1

       

      I don't know if that it's supposed to be supported by OpenCL language but

      some samples that work from NVIDIA OpenCL SDK doesn't work with your SDK mainly is because they use something like:

      __kernel void filter( __global const uchar4* src, __global unsigned int* dest,

      __local uchar4 local[32][32])

       

      The problem lies in __local uchar4 local[32][32] your SDK fails to compile..

      A  solution I have found to fix  their examples for your SDK is to change two things:

      1. __local uchar4 local[32][32] by __local *local in kernel definiton

      and

      2. References to the array

      change local[x][y] with local[y+(32)*x]

      Are you going to fix this or are NVIDIA guys using non standard features?

       

       

       

       

        • 2 dimensional fixed size arrays supported?
          brg

          OpenCL allows for the static allocation of local array sizes only within a kernel and not as arguments. This may change in the future and indeed we may choose to implement this but it is not a supported feature, today, either in the specification or our implementation.

          As you have pointed out according to the specification the standard way of passing local arugments is through arguments of the form:

             __local T * name

          This is defined in the OpenCL 1.0 specification as:

          "6.5.2 __local (or local)
          The __local or local address space name is used to describe variables that need to be allocated in local memory and are shared by all work-items of a work-group. This qualifier can be used with arguments to functions (including __kernel functions) declared as pointers, or with variables declared inside a __kernel function."

          Note the sentance that begins with "This qualifier can be used with arguments to functions (including __kernel functions) declared as pointers,"

            • 2 dimensional fixed size arrays supported?
              ochafik

              Hi,

              Sorry for respawning this month-and-a-half old topic...

               

              Originally posted by: brg OpenCL allows for the static allocation of local array sizes only within a kernel and not as arguments.


              I believe there is no static allocation taking place inside the function here.

              The syntax "void f(int x[64])" does not allocate an int[64] inside the function f. It just types the pointer x as being a pointer to an int[64] array that must be allocated by the caller, the outside array is indeed not passed by value (not to be confused with non-argument variable declaration).

              As the local arguments must be allocated through the API by hand anyway, can't see much of an issue supporting this syntax here...

               

               

              "6.5.2 __local (or local) The __local or local address space name is used to describe variables that need to be allocated in local memory and are shared by all work-items of a work-group. This qualifier can be used with arguments to functions (including __kernel functions) declared as pointers, or with variables declared inside a __kernel function."

              Note the sentance that begins with "This qualifier can be used with arguments to functions (including __kernel functions) declared as pointers,"



              Hence I'd say it is technically our case here, if we agree that a non-allocated-array that behaves as a pointer is assimilable to a pointer.

              Otherwise, good job on ATI Stream !

              Regards

                • 2 dimensional fixed size arrays supported?
                  jcpalmer

                  I saw a post on the NVIDIA board, where someone pointed out that this was not proper, see OpenCL Specification 5.5.2, pg 103.  "If the argument is declared with the __local qualifier, the arg_value entry must be NULL".  Someone from NVIDIA agreed at the end.

                  Hope this is not improper, but it was

                  http://forums.nvidia.com/index.php?showtopic=104785

                    • 2 dimensional fixed size arrays supported?
                      ochafik

                      @jcpalmer Thanks for the link

                       

                      "If the argument is declared with the __local qualifier, the arg_value entry must be NULL"


                      This makes sense, as there is not one local variable but one local variable for each local scope. The size given to clSetKernelArg will be used to allocate the local values in each scope properly.

                       

                      Someone from NVIDIA agreed at the end.


                      Only makes it more likely ;-)

                      What appals me with a "__local *local" syntax is that it means using a void* where we could use proper typing.

                      If local pointers can be properly typed (the spec gives an example with "__local half *pl"), why couldn't local int array pointers be too, as they are technically equivalent ?