avinashcpandey/ocl-radix-sort

Automatically exported from code.google.com/p/ocl-radix-sort

ocl-radix-sort's Issues

Not enough local memory

Dear Philippe Helluy.

Congrats for your amazing work!

I'm Jose Luis Cercós Pita (another time) with a small bug.

If you select _BITS = 8 (CLRadixSortParam.hpp), local memory is not enough into 
some NVidia cards, but any OUT OF MEMORY is reported. In this this case the 
kernels simply pass without any operations, with a broken result.

To solve it (or test it), you can add something like this:

    cl_ulong localMem;
    clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(localMem), &localMem, NULL);
    if(localMem < sizeof(cl_uint)*_RADIX*_ITEMS){
        ERROR
    }
    unsigned int maxmemcache=max(_HISTOSPLIT,_ITEMS * __RADIX_GROUPS__ * _RADIX / _HISTOSPLIT);
    if(localMem < sizeof(cl_uint)*maxmemcache){
    ERROR
    }

Regards,
Jose Luis Cercós Pita

Original issue reported on code.google.com by [email protected] on 30 May 2011 at 2:09

Not working on OSX Radeon HD 4850

SVN revision 52 after compiling:

crb002$ ./go

Test CLRadixSort class...

Available platforms:
Platform 1:
Apple
Apple
OpenCL 1.0 (Dec 23 2010 17:30:26)
We choose Device 1/2 of Platform 1/1:
Radeon HD 4850
OpenCL Extensions for that device: cl_APPLE_gl_sharing

Sort on GPU
GPU cache=16384
Needed cache=8192
Compute units=10
Create the context
Create the command queue
OpenCL initializations OK !

Cache size=16384 Bytes
Needed cache=8192 Bytes
Construct the random list
Send to the GPU
Resize to  1048576
Radix=32
Max Int=536870912
sorting 1048576 keys

Start storting 1048576 keys
Transpose
pass 0
Build histograms 
Scan histograms 
[CL_INVALID_WORK_GROUP_SIZE] : OpenCL Error : clEnqueueNDRangeKernel failed: 
total work group size (256) is greater than the device can support (64)
Break on OpenCLErrorBreak to debug.
Assertion failed: (err== CL_SUCCESS), function ScanHistogram, file 
CLRadixSort.cpp, line 789.
Abort trap

Original issue reported on code.google.com by [email protected] on 25 Jul 2011 at 3:57

Changes for sorting 64-bit keys

What all changes need to be incorporated so as to make it sort 64 bit keys 
(cl_ulong type keys). Please guide.

Original issue reported on code.google.com by [email protected] on 28 Jan 2015 at 6:31

Variable start never referenced

If you define TRANSPOSE, variable start is created and evaluated, but never 
used.

I suggest something like this:

#ifndef TRANSPOSE
  unsigned int start= ig *(n/groups/items);
#endif

Regards,
Jose Luis Cercós Pita

Original issue reported on code.google.com by [email protected] on 21 Jun 2011 at 9:08

Makefile (tested in Debian Linux)

Hi Philippe Helluy.

I built a makefile tested into Ubuntu/Debian Linux distributions.

I hope can help you!

Regards,
Jose Luis Cercós Pita

Original issue reported on code.google.com by [email protected] on 16 Jun 2011 at 1:30

Attachments:

• makefile

License on source codes

CLRadixSort.cl hasn't license. It is a little bit dangeorus.

Regards,
Jose Luis Cercós Pita

Original issue reported on code.google.com by [email protected] on 16 Jun 2011 at 1:42

Key Value sort

Can this library sort Key-Value pairs as well?

Original issue reported on code.google.com by [email protected] on 3 Aug 2011 at 4:10

Max unsigned integer trouble

Hi! My name is Jose Luis Cercós Pita, and i will implement your radix sort 
implementation in my code.

I'm using 0-2 version.

The problem is associated to _MAXINT defined variable (CLRadixSortParam.hpp). 
If the unsigned integer type selected has 32 bits (#define _TOTALBITS 32) 
_MAXINT overflows the maximum capability of the unsigned int type (because 2^0 
has not taken into account).

In order to solve it, I suggest you 3 changes:

CLRadixSortParam.hpp (line 31):
[-] #define _MAXINT (1 << _TOTALBITS)
[+] #define _MAXINT (1 << (_TOTALBITS-1))

CLRadixSort.cpp (line 225):
[-] int pad[_GROUPS * _ITEMS];
[+] unsigned int pad[_GROUPS * _ITEMS];

CLRadixSort.cpp (line 227):
[-] pad[ii]=_MAXINT - 1;
[+] pad[ii]=_MAXINT - (unsigned int)1;

I tested it, and runs OK!
I hope can help!

Regards,
Jose Luis Cercós Pita

Original issue reported on code.google.com by [email protected] on 19 May 2011 at 8:23