I am writing cuda version of merge sort and if I am using cudaMemcpyDeviceToHost in order to get back list of elements from GPU, it's giving memory error, on the other side if I am commenting out the line then the program is not sorting properly. Can anyone please suggest.
/* C program for Merge Sort with Cuda Technology*/
#include<stdlib.h>
#include<stdio.h>
#include <cuda.h>
#include <sys/time.h>
#define THR1 1000
#define THR2 10000
#define N 800000
/*
********************************
Program UTILITY Code here
********************************
*/
static void HandleError( cudaError_t err, const char *file, int line ) {
if (err != cudaSuccess) {
printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
exit( EXIT_FAILURE );
}}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
/* Function to print an array */
void printArray(int A[], int size)
{
int i;
for (i=0; i < size; i++)
printf("%d ", A[i]);
printf("\n");
}
//Function to test if the output is in asending order or not
void test(int a[], int n) {
int i;
for (i=1;i<n;++i) {
if (a[i]<a[i-1]) {
break;
}
}
if (i<n) {
for (i=1;i<n;++i) {
if (a[i]>a[i-1]){
break;
}
}
if (i<n) {
printf("\nArray is not sorted\n");
}
}
else {
printf("\nArray is sorted\n");
}
}
/*
*****************************************
Sequential Version here
*****************************************
*/
void insertionSort(int array[], int min, int max)
{
int key ;
// we loop through all elements in the original array from the min + 1 element
for (int j = min + 1 ; j <= max ; j++)
{
// store the current element as the key
key = array[j] ;
// get the element just before the current element
int i = j - 1 ;
// loop through all elements from the key to the min element
// check if the current element is smaller than the key
while (i >= min && array[i] > key)
{
// we move the current element backward
array[i+1] = array[i] ;
i-- ;
}
// we finally move the key
array[i+1] = key ;
}
}
void merge(int array[], int min, int max, int mid)
{
int firstIndex = min;
int secondIndex = mid + 1;
int * tempArray = new int [max + 1];
// While there are elements in the left or right runs
for (int index = min; index <= max; index++) {
// If left run head exists and is <= existing right run head.
if (firstIndex <= mid && (secondIndex > max || array[firstIndex] <= array[secondIndex]))
{
tempArray[index] = array[firstIndex];
firstIndex = firstIndex + 1;
}
else
{
tempArray[index] = array[secondIndex];
secondIndex = secondIndex + 1;
}
}
// transfer to the initial array
for (int index = min ; index <= max ; index++)
array[index] = tempArray[index];
}
void smergeSort(int array[], int min, int max, int threshold)
{
// prerequisite
if ( (max - min + 1) <= threshold )
{
insertionSort(array, min, max);
}
else
{
// get the middle point
int mid = (max+min) / 2;
// apply merge sort to both parts of this
smergeSort(array, min, mid, threshold);
smergeSort(array, mid+1, max, threshold);
// and finally merge all that sorted stuff
merge(array, min, max, mid) ;
}
}
/*
*****************************************
Parallel Version here
*****************************************
*/
__device__ void gpu_bottomUpMerge(int* source, int* dest, int start, int middle, int end) {
int i = start;
int j = middle;
for (int k = start; k < end; k++) {
if (i < middle && (j >= end || source[i] < source[j])) {
dest[k] = source[i];
i++;
} else {
dest[k] = source[j];
j++;
}
}
}
__global__ void gpu_mergesort(int* source, int* dest, int size, int width, int slices, dim3* threads, dim3* blocks) {
int idx = blockDim .x * blockIdx .x + threadIdx .x;
int start = width*idx*slices,
middle,
end;
for (int slice = 0; slice < slices; slice++) {
if (start >= size)
break;
middle = min(start + (width >> 1), size);
end = min(start + width, size);
gpu_bottomUpMerge(source, dest, start, middle, end);
start += width;
}
}
void mergesort(int* data, int size, dim3 threadsPerBlock, dim3 blocksPerGrid) {
// Allocate two arrays on the GPU we switch back and forth between them during the sort
int* D_data;
int* D_swp;
dim3* D_threads;
dim3* D_blocks;
// Actually allocate the two arrays
HANDLE_ERROR(cudaMalloc((void**) &D_data, size * sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**) &D_swp, size * sizeof(int)));
// Copy from our input list into the first array
HANDLE_ERROR(cudaMemcpy(D_data, data, size * sizeof(int), cudaMemcpyHostToDevice));
int* A = D_data;
int* B = D_swp;
int nThreads = threadsPerBlock.x * threadsPerBlock.y * threadsPerBlock.z * blocksPerGrid.x * blocksPerGrid.y * blocksPerGrid.z;
// Divide the list and give pieces of it to each thread, letting the pieces grow bigger and bigger until the whole list is sorted
for (int width = 2; width < (size << 1); width <<= 1) {
int slices = size / ((nThreads) * width) + 1;
// Actually call the kernel
gpu_mergesort<<<blocksPerGrid, threadsPerBlock>>>(A, B, size, width, slices, D_threads, D_blocks);
cudaDeviceSynchronize();
// Switch the input / output arrays instead of copying them around
A = A == D_data ? D_swp : D_data;
B = B == D_data ? D_swp : D_data;
}
// Get the list back from the GPU
HANDLE_ERROR(cudaMemcpy(data, A, size * sizeof(int), cudaMemcpyDeviceToHost));
// Free the GPU memory
HANDLE_ERROR(cudaFree(A));
HANDLE_ERROR(cudaFree(B));
}
/* Driver program to test above functions */
int main()
{
dim3 threadsPerBlock;
dim3 blocksPerGrid;
threadsPerBlock.x = 224;
blocksPerGrid.x = 10;
int i, *a;
printf("How many elements in the array? ");
a = (int *)malloc(sizeof(int) * N);
srand(time(0));
for(i=0;i<N;i++)
{
a[i]=rand()%1000;
}
printf("List Before Sorting...\n");
// printArray(a, N);
if (N<=THR2)
{
clock_t begin = clock();
smergeSort(a, 0, N - 1, THR2);
clock_t end = clock();
printf("\nSorted array: ");
//printArray(a,N);
printf("\n");
test(a,N);
double time_spent = (double)(end - begin) / CLOCKS_PER_SEC;
printf("SM");
printf("Elapsed: %f seconds\n",time_spent );
printf("\nSize of the array is %d",N);
exit(0);
}
else
{
clock_t begin = clock();
mergesort(a, N, threadsPerBlock, blocksPerGrid);
clock_t end = clock();
printf("\nSorted array: ");
//printArray(a,N);
printf("\n");
test(a,N);
double time_spent = (double)(end - begin) / CLOCKS_PER_SEC;
printf("Cuda");
printf("Elapsed: %f seconds\n",time_spent );
printf("\nSize of the array is %d\n",N);
exit(0);
}
}
Now program is working fine even for large elements, however, when I am using a large number of threads, let's say block 10 and threads 224 it;'s giving error :- an illegal memory access was encountered in mergesort.cu at line 215
After debugging the code I am getting below errors again:-
========= Invalid __global__ read of size 4
========= at 0x00000148 in
/home/sharmpra/mergesort.cu:150:gpu_mergesort(int*, int*, int, int, int, dim3*, dim3*)
========= by thread (96,0,0) in block (9,0,0)
========= Address 0x915fc0000 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/lib64/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x204235]
========= Host Frame:./a.out [0x1e831]
========= Host Frame:./a.out [0x3c3d3]
========= Host Frame:./a.out [0x38a8]
========= Host Frame:./a.out [0x37b1]
========= Host Frame:./a.out [0x3810]
========= Host Frame:./a.out [0x33d1]
========= Host Frame:./a.out [0x35ae]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf0) [0x20790]
========= Host Frame:./a.out [0x2bc9]
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/lib64/libcuda.so.1 [0x2ef503]
========= Host Frame:./a.out [0x3c0f6]
========= Host Frame:./a.out [0x33da]
========= Host Frame:./a.out [0x35ae]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf0) [0x20790]
========= Host Frame:./a.out [0x2bc9]
=========
May please someone suggest what additional things I can implement in the program to remove such errors. Also, I am using this command line settings:-nvcc -o a.out -Wno-deprecated-gpu-targets -lineinfo -arch=compute_20,sm_20 -rdc=true -lcudadevrt mergesort.cu
