How to set bindings?

[Dinh]

Hi,

Everything was going great when I was setting up my variables to pass
to the kernel. Now, as I try to pass these variables, I get this
exception:

>>terminate called after throwing an instance of 'cupp::exception::kernel_parameter_type_mismatch'
>> what(): Wrong parameter type passed to kernel. Maybe forgot to set bindings?

Please let me know what's going on!! Thanks!

My kernel call is:
RunAlignKernel(d,
d_indexN1,
d_indexId1,
d_indexLoc1,
d_bseq,
d_reg,
d_seeds0,
d_profile0a,
d_profile0b1,
d_readlen);

My kernel_t:

typedef void(*kernelT)(cupp::deviceT::vector< unsigned > &,
cupp::deviceT::vector< cupp::deviceT::vector< unsigned > > &,
cupp::deviceT::vector< cupp::deviceT::vector< unsigned > > &,
cupp::deviceT::vector< cupp::deviceT::vector< bit24_td > >&,
cupp::deviceT::vector< cupp::deviceT::vector< bit24_td > >&,
cupp::deviceT::vector< uint4 > &,
cupp::deviceT::vector< short4 > &,
cupp::deviceT::vector< short4 > &,
cupp::deviceT::vector< unsigned > &);

And finally, my kernel set up:

__global__ void global_function (deviceT::vector< unsigned > &d_index,
deviceT::vector< deviceT::vector< unsigned > > &d_indexId,
deviceT::vector< deviceT::vector< unsigned > > &d_indexLoc,
deviceT::vector< deviceT::vector< bit24_td > >& d_bseq,
deviceT::vector< deviceT::vector< bit24_td > >& d_reg,
deviceT::vector< uint4 > &d_seed0,
deviceT::vector< short4 > &d_profile0a,
deviceT::vector< short4 > &d_profile0b1,
deviceT::vector< unsigned > &d_readlen)

[Jens]

Sorry, there was a bug when using the cuda vector types. I'll send you
a fix by mail, as I can't attach a file here. If that won't help, can
you please post the type of the variables you pass to the kernel?

Type bindings are set up by defining for example

struct test {
typedef deviceT::test device_type;
typedef test host_type;
};

in a struct/class.

-Jens

BTW: In case you don't change the elements of some vectors you pass to
the kernel, I would suggest defining them as 'const' so CuPP will not
transfer the data back from the GPU.

[Dinh Diep]

Hello Jens,

I do not want to get these value returned but I am copying values from a large data set into these smaller vectors to pass to the kernel. If I define them as const, how would I be able to write to them?

I've been avoiding the need to pass struct/class since I still don't know how to do the type transformation. But it seems like I will have to.

The types which I passed were:

    cupp::vector <unsigned> d_indexN1;
    cupp::vector < cupp::vector<unsigned> > d_indexId1;
    cupp::vector < cupp::vector<unsigned> > d_indexLoc1;
cupp::vector < cupp::vector<bit24_t> > d_bseq;
cupp::vector < cupp::vector<bit24_t> > d_reg;
    cupp::vector<uint4> d_seeds0;
    cupp::vector<short4> d_profile0a;
    cupp::vector<short4> d_profile0b1;
    cupp::vector<short> d_readlen;

bit24_t is a structure:

struct bit24_t
{
unsigned a:24;

}

I've tried to set the binding as follow:

struct bit24_td;
struct bit24_t
{
    unsigned a:24;
    typedef bit24_td device_type;
    typedef bit24_t host_type;
};

struct bit24_td
{
    unsigned a:24;
    typedef bit24_td device_type;
    typedef bit24_t host_type;
};

This must have been incorrect as I got a long list of compile errors, all relating to bit24_t.

-Dinh

[Jens]

Hi,

On 3 Mrz., 19:28, Dinh Diep <hdin...@gmail.com> wrote:
> I do not want to get these value returned but I am copying values from a
> large data set into these smaller vectors to pass to the kernel. If I define
> them as const, how would I be able to write to them?

Sorry. Of course they need to be non-const. However, if only read
values inside a kernel, you should define them as const.

>
> I've been avoiding the need to pass struct/class since I still don't know
> how to do the type transformation. But it seems like I will have to.

Actually, it should be possible to directly pass struct bit24_t to a
kernel as it is C complaint. Type transformations are only necessary
if you have a C++ class/structure or you want to use a different kind
of data representation on the device. Is directly passing bit24_t to a
kernel fixing your problem? In case not, I would suggest that you
outcomment arguments, to see when the error disappears. As far as I
can see, it should work fine.

-Jens

[Dinh]

Indeed I was able to pass the struct object by itself. However, I
really need to pass a 2D array of such objects. I had tried to pass
them inside of vectors but it was unsuccessful. Is there a better way
to pass a bit24_t [N][N] object?

[Jens]

Have you tried passing a

cupp::vector < cupp::vector<bit24_t> >

to the kernel and expect

cupp::deviceT::vector < cupp::deviceT::vector<bit24_t> >

? If you can pass a bit24_t this should work just fine. I case it is
not working, please post the error message (or send me a copy of the
code so I can compile it here)

On Mar 4, 6:19 am, Dinh <hdin...@gmail.com> wrote:
> Is there a better way
> to pass a bit24_t [N][N] object?

This depends on what you need. I would use vectors, but you may as
well just use a memory1d object and manually do the 2-dimensional
addressing. Or if you don't need any of the additional functionality
just use CUDA malloc/memcpy functions (or their CuPP counterparts).

-Jens