I have made the following code to be able to be used to sort an array of char * using THRUST sort method. For some reason the code is breaking whenever it tries to compare the chars in the string.
thrust::sort_by_key(deviceArrayToSort.begin(),deviceArrayToSort.end(),deviceArrayToSortRow.begin(),CharArrayCmp());
The comparator is below:
struct CharArrayCmp{
__host__ __device__
bool operator()(const CharArr & o1, const CharArr & o2) {
return this->compare(o1.value,o1.length,o2.value,o2.length);
}
__host__ __device__ bool compare (const char * src, int lenSrc, const char * dst, int lenDest)
{
int end;
if(lenSrc > lenDest){
end = lenDest;
}else{
end = lenSrc;
}
for(int i = 0; i < end; i++){
if(src[i] > dst[i]){
return false;
}
}
if(lenSrc >= lenDest){
return false;
}
return true;
}
};
When it tries to run on the device on this line there is an error:
if(src[i] != dst[i])
thrust::system::cuda::detail::detail::launch_closure_by_value<thrust::system::cuda::detail::
detail::stable_sort_by_count_detail::stable_sort_by_count_closure<256u,
thrust::detail::normal_iterator<thrust::pointer<unsigned int,
thrust::system::cuda::detail::tag, thrust::use_default, thrust::use_default> >,
thrust::detail::normal_iterator<thrust::pointer<unsigned int,
thrust::system::cuda::detail::tag, thrust::use_default, thrust::use_default> >,
thrust::system::cuda::detail::temporary_indirect_ordering<thrust::system::cuda::detail::tag,
thrust::detail::normal_iterator<thrust::device_ptr<CharArr> >, CharArrayCmp>::compare,
thrust::system::cuda::detail::detail::statically_blocked_thread_array<256u> > > [0] [device
0 (GK104)] (Signal) - CUDA_EXCEPTION_10:Device Illegal Address
CUDA Thread (193,0,0) Block (22,0,0)
All Kernel Threads (144 Blocks of 256 Threads)
I am quite new to CUDA so I am uncertain of what I am doing wrong but it feels like this should be pretty straightforward.
This is the CharArr struct:
typedef struct{
char * value;
int length;
} CharArr;
And finally here is the code that makes use of this sort_by_key. I have made sure that the information being passed into this is correct. Namely that arrayToSort and arrayToSortRow are both arrays, of char * and long long respectively and size is the size of these two arrays.
void sortCharArrayStable(char ** arrayToSort, long long * arrayToSortRow,long long size){
std::cout <<"about to start LongIndex" <<std::endl;
thrust::host_vector<CharArr> hostToSort(size);
thrust::host_vector<long long> hostToSortRow(size);
for(int i =0; i < size; i++){
CharArr sortRow;
if(arrayToSort[i] == 0x0){
sortRow.length = 0;
sortRow.value = "";
std::cout<<"Had an error on row "<< arrayToSortRow[i]<<" when making column array for sortCharArrayStable"<<std::endl;
}else{
sortRow.length = strlen(arrayToSort[i]);
sortRow.value = arrayToSort[i];
}
hostToSort[i] = sortRow;
hostToSortRow[i] = arrayToSortRow[i];
}
thrust::device_vector<CharArr> deviceArrayToSort = hostToSort;// (arrayToSort,arrayToSort + size);
thrust::device_vector<long long> deviceArrayToSortRow = hostToSortRow;
thrust::stable_sort_by_key(deviceArrayToSort.begin(),deviceArrayToSort.end(),deviceArrayToSortRow.begin(),CharArrayCmp());
//copy the contents back into our original array to sort now sorted
hostToSort = deviceArrayToSort;
for(int i =0; i < size; i++){
arrayToSort[i] = hostToSort[i].value;
}
arrayToSortRow);
thrust::copy(deviceArrayToSortRow.begin(),deviceArrayToSortRow.end(),arrayToSortRow);
}
Here is a complete example of the issue ocurring in a compilable example:
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/reduce.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
typedef struct{
char * value;
int length;
} CharArr;
struct CharArrayCmp{
__host__ __device__
bool operator()(const CharArr & o1, const CharArr & o2) {
return this->compare(o1.value,o1.length,o2.value,o2.length);
}
__host__ __device__ bool compare (const char * src, int lenSrc, const char * dst, int lenDest)
{
int end;
if(lenSrc > lenDest){
end = lenDest;
}else{
end = lenSrc;
}
for(int i = 0; i < end; i++){
if(src[i] > dst[i]){
return false;
}
}
if(lenSrc >= lenDest){
return false;
}
return true;
}
};
void sortCharArray(char ** arrayToSort, long long * arrayToSortRow,long long size){
std::cout <<"about to start LongIndex" <<std::endl;
thrust::host_vector<CharArr> hostToSort(size);
thrust::host_vector<long long> hostToSortRow(size);
for(int i =0; i < size; i++){
CharArr sortRow;
sortRow.value = arrayToSort[i];
sortRow.length = strlen(arrayToSort[i]);
hostToSort[i] = sortRow;
hostToSortRow[i] = arrayToSortRow[i];
}
thrust::device_vector<CharArr> deviceArrayToSort = hostToSort;// (arrayToSort,arrayToSort + size);
thrust::device_vector<long long> deviceArrayToSortRow = hostToSortRow;//(arrayToSortRow,arrayToSortRow + size);
// thrust::sort(deviceArrayToSort.begin(),deviceArrayToSort.end());
thrust::sort_by_key(deviceArrayToSort.begin(),deviceArrayToSort.end(),deviceArrayToSortRow.begin(),CharArrayCmp());
//copy the contents back into our original array to sort now sorted
hostToSort = deviceArrayToSort;
for(int i =0; i < size; i++){
arrayToSort[i] = hostToSort[i].value;
}
thrust::copy(deviceArrayToSortRow.begin(),deviceArrayToSortRow.end(),arrayToSortRow);
}
int main()
{
char ** charArr = new char*[10];
charArr[0] = "zyxw";
charArr[1] = "abcd";
charArr[2] = "defg";
charArr[3] = "werd";
charArr[4] = "aasd";
charArr[5] = "zwedew";
charArr[6] = "asde";
charArr[7] = "rurt";
charArr[8] = "ntddwe";
charArr[9] = "erbfde";
long long * rows = new long long[10];
for(int i = 0; i < 10;i++ ){
rows[i] = i;
}
sortCharArray(charArr,rows,10);
for(int i = 0; i < 10; i++){
std::cout<<"Row is "<<rows[i]<<" String is "<<charArr[i]<<std::endl;
}
}
This will not work:
thrust::device_vector<CharArr> deviceArrayToSort = hostToSort;// (arrayToSort,arrayToSort + size);
thrust::device_vector<long long> deviceArrayToSortRow = hostToSortRow;
thrust::stable_sort_by_key(deviceArrayToSort.begin(),deviceArrayToSort.end(),deviceArrayToSortRow.begin(),CharArrayCmp());
Each CharArr object in hostToSort contained a pointer that points to a host memory location. The numerical value of this pointer is unchanged when you copy it to a device vector. If you then try and dereference this (bogus) pointer in device code, it will fail exactly as you are seeing.
You will need to go through the deviceArrayToSort device vector, and for each CharArr object in it, you will need to adjust its value pointer to point to a valid location in device memory, which would presumably be the starting address of each character string to sort.