Results 1 to 9 of 9

Thread: how to declare a constant array of float2 vectors?

  1. #1
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    how to declare a constant array of float2 vectors?

    In my kernel code I am trying to declare a constant array of float2 vectors like this:
    Code :
    __constant float2 grads[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};

    But this gives a rather obscure error when attempting to build the kernel:

    ptxas application ptx input, line 11; fatal : Parsing error near ',': syntax error
    ptxas fatal : Ptx assembly aborted due to errors
    error : Ptx compilation failed: gpu='sm_20', device code='cuModuleLoadDataEx_4'
    : Considering profile 'compute_20' for gpu='sm_20' in 'cuModuleLoadDataEx_4'
    : Retrieving binary for 'cuModuleLoadDataEx_4', for gpu='sm_20', usage mode=' '
    : Considering profile 'compute_20' for gpu='sm_20' in 'cuModuleLoadDataEx_4'
    : Control flags for 'cuModuleLoadDataEx_4' disable search path
    : Ptx binary found for 'cuModuleLoadDataEx_4', architecture='compute_20'
    : Ptx compilation for 'cuModuleLoadDataEx_4', for gpu='sm_20', ocg options=' '

    Any suggestions for the correct way to do this?

    I'm using the Linux Nvidia OpenCL implementation.
    This code did previously work ok on Mac, so I was thinking it was correct syntax...

    Thanks!

  2. #2
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: how to declare a constant array of float2 vectors?

    Did you declare grads at program scope or at function scope? That is, does it look like "foo" or like "bar" in the code below:

    Code :
    // Program scope. This is valid and should compile.
    __constant float2 foo[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};
     
    __kernel void fubar()
    {
        // Function scope. This is not valid and should not compile.
        __constant float2 foo[2] = {(float2)(1.0f,0.0f), (float2)(0.0f,1.0f)};
    }

    Section 6.5.3. of the spec says:
    Variables allocated in the __constant address space can only be defined as program scope variables and are required to be initialized.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  3. #3
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    Re: how to declare a constant array of float2 vectors?

    Yes it is at program scope. (not function scope)

  4. #4
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    Re: how to declare a constant array of float2 vectors?

    ...so it should be correct syntax and should work, but doesn't. Any ideas? Bug?

  5. #5
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: how to declare a constant array of float2 vectors?

    Yeah, as I mentioned earlier, if it is in program scope it should work. There's a compiler bug.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  6. #6

    Re: how to declare a constant array of float2 vectors?

    Thanks - had the same problem now it's resolved!

  7. #7
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    Re: how to declare a constant array of float2 vectors?

    Care to share with us how it got resolved?

  8. #8
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    Re: how to declare a constant array of float2 vectors?

    just an update on this, in case anybody else is running into the same issue:
    Nvidia says it should be fixed in their driver r280 release.

  9. #9
    Junior Member
    Join Date
    Aug 2009
    Posts
    9

    Re: how to declare a constant array of float2 vectors?

    Quote Originally Posted by ksi
    just an update on this, in case anybody else is running into the same issue:
    Nvidia says it should be fixed in their driver r280 release.
    I did run into the same issue (with driver 306.94) and did some further investigations, see https://devtalk.nvidia.com/default/topic/524095.

    Kind regards,
    Markus

Similar Threads

  1. Casting int2, float2
    By toneburst in forum OpenCL
    Replies: 1
    Last Post: 12-04-2010, 06:49 PM
  2. Replies: 3
    Last Post: 07-04-2010, 02:57 AM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •