0

I have a CUDA kernel which takes an edge image and processes it to create a smaller, 1D array of the edge pixels. Now here is the strange behaviour. Every time I run the kernel and calculate the number of edge pixels in "d_nlist" (see the code near the printf), I get a greater pixel count each time, even when I use the same image and stop the program completely and re-run. Therefore, each time I run it, it takes longer to run, until eventually, it throws an un-caught exception.

My question is, how can I stop this from happening so that I can get consistent results each time I run the kernel?

My device is a Geforce 620.

Constants:

THREADS_X = 32
THREADS_Y = 4
PIXELS_PER_THREAD = 4
MAX_QUEUE_LENGTH = THREADS_X * THREADS_Y * PIXELS_PER_THREAD
IMG_WIDTH = 256
IMG_HEIGHT = 256
IMG_SIZE = IMG_WIDTH * IMG_HEIGHT
BLOCKS_X = IMG_WIDTH / (THREADS_X * PIXELS_PER_THREAD)
BLOCKS_Y = IMG_HEIGHT / THREADS_Y

The kernel is as follows:

__global__ void convert2DEdgeImageTo1DArray( unsigned char const * const image, 
unsigned int* const list, int* const glob_index ) {

unsigned int const x = blockIdx.x  * THREADS_X*PIXELS_PER_THREAD + threadIdx.x;
unsigned int const y = blockIdx.y  * THREADS_Y + threadIdx.y;

volatile int qindex = -1;
volatile __shared__ int sh_qindex[THREADS_Y];
volatile __shared__ int sh_qstart[THREADS_Y];
sh_qindex[threadIdx.y] = -1;

// Start by making an array
volatile __shared__ unsigned int sh_queue[MAX_QUEUE_LENGTH];

// Fill the queue
for(int i=0; i<PIXELS_PER_THREAD; i++)
{
    int const xx = i*THREADS_X + x;

    // Read one image pixel from global memory
    unsigned char const pixel = image[y*IMG_WIDTH + xx];
    unsigned int  const queue_val = (y << 16) + xx;

    if(pixel)
    {           
        do {
            qindex++;
            sh_qindex[threadIdx.y] = qindex;
            sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + qindex] = queue_val;
        } while (sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + qindex] != queue_val);
    }

    // Reload index from smem (last thread to write to smem will have updated it)
    qindex = sh_qindex[threadIdx.y];
}

// Let thread 0 reserve the space required in the global list
__syncthreads();
if(threadIdx.x == 0 && threadIdx.y == 0)
{
    // Find how many items are stored in each list
    int total_index = 0;
    #pragma unroll
    for(int i=0; i<THREADS_Y; i++)
    {
        sh_qstart[i] = total_index;
        total_index += (sh_qindex[i] + 1u);
    }

    // Calculate the offset in the global list
    unsigned int global_offset = atomicAdd(glob_index, total_index);
    #pragma unroll
    for(int i=0; i<THREADS_Y; i++)
    {
        sh_qstart[i] += global_offset;
    }
}
__syncthreads();

// Copy local queues to global queue
for(int i=0; i<=qindex; i+=THREADS_X)
{
    if(i + threadIdx.x > qindex)
        break;

    unsigned int qvalue = sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + i + threadIdx.x];
    list[sh_qstart[threadIdx.y] + i + threadIdx.x] = qvalue;
}
}

The following is the method which calls the kernel:

void call2DTo1DKernel(unsigned char const * const h_image)
{
    // Device side allocation
    unsigned char *d_image = NULL;
    unsigned int *d_list = NULL;
    int h_nlist, *d_nlist = NULL;
    cudaMalloc((void**)&d_image, sizeof(unsigned char)*IMG_SIZE);
    cudaMalloc((void**)&d_list, sizeof(unsigned int)*IMG_SIZE);
    cudaMalloc((void**)&d_nlist, sizeof(int));

    // Time measurement initialization
    cudaEvent_t start, stop, startio, stopio;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventCreate(&startio); 
    cudaEventCreate(&stopio);

    // Start timer w/ io
    cudaEventRecord(startio,0);

    // Copy image data to device
    cudaMemcpy((void*)d_image, (void*)h_image, sizeof(unsigned char)*IMG_SIZE,    cudaMemcpyHostToDevice);

    // Start timer
    cudaEventRecord(start,0);

    // Kernel call
    // Phase 1 : Convert 2D binary image to 1D pixel array
    dim3 dimBlock1(THREADS_X, THREADS_Y);
    dim3 dimGrid1(BLOCKS_X, BLOCKS_Y);
    convert2DEdgeImageTo1DArray<<<dimGrid1, dimBlock1>>>(d_image, d_list, d_nlist);

    // Stop timer
    cudaEventRecord(stop,0);
    cudaEventSynchronize(stop);

    // Stop timer w/ io
    cudaEventRecord(stopio,0);
    cudaEventSynchronize(stopio);

    // Time measurement
    cudaEventElapsedTime(&et,start,stop);
    cudaEventElapsedTime(&etio,startio,stopio);

    // Time measurement deinitialization
    cudaEventDestroy(start); 
    cudaEventDestroy(stop);
    cudaEventDestroy(startio); 
    cudaEventDestroy(stopio);

    // Get list size
    cudaMemcpy((void*)&h_nlist, (void*)d_nlist, sizeof(int), cudaMemcpyDeviceToHost);

    // Report on console
    printf("%d pixels processed...\n", h_nlist);

    // Device side dealloc
    cudaFree(d_image);
    cudaFree(d_space);
    cudaFree(d_list);
    cudaFree(d_nlist);
}

Thank you very much in advance for your help everyone.

Adam
  • 610
  • 1
  • 7
  • 21
  • Do you have an actual question here? All I see is a lot of code and a list symptoms. What *exactly* is wrong and, equally as important, why is it wrong? What are you expecting an answer will tell you? Help us help you.... – talonmies Aug 07 '13 at 11:00
  • Well basically the number of pixels processed should be the same each time as I am using the same image. The problem is, the printf gives me a different result every time. The program should literally read in the image, and put the edge pixels into a list, which is smaller than the original image. Therefore, I should get the same number of pixels in the array each time. My question is I guess, how can I stop this from happening? Hope this helps. – Adam Aug 07 '13 at 11:09
  • So where is the `houghKernel2_3_phase1` in the code you have shown? – sgarizvi Aug 07 '13 at 15:01
  • Sorry, my mistake, I've renamed the method to convert2DEdgeImageTo1DArray. – Adam Aug 07 '13 at 15:11

1 Answers1

1

As a preamble, let me suggest some troubleshooting steps that are useful:

  1. instrument your code with proper cuda error checking
  2. run your code with cuda-memcheck e.g. cuda-memcheck ./myapp

If you do the above steps, you'll find that your kernel is failing, and the failures have to do with global writes of size 4. So that will focus your attention on the last segment of your kernel, beginning with the comment // Copy local queues to global queue

Regarding your code, then, you have at least 2 problems:

  1. The addressing/indexing in your final segment of your kernel, where you are writing the individual queues out to global memory, is messed up. I'm not going to try and debug this for you.
  2. You are not initializing your d_nlist variable to zero. Therefore when you do an atomic add to it, you are adding your values to a junk value, which will tend to increase as you repeat the process.

Here's some code which has the problems removed, (I did not try to sort out your queue copy code) and error checking added. It produces repeatable results for me:

$ cat t216.cu
#include <stdio.h>
#include <stdlib.h>

#define THREADS_X 32
#define THREADS_Y 4
#define PIXELS_PER_THREAD 4
#define MAX_QUEUE_LENGTH (THREADS_X*THREADS_Y*PIXELS_PER_THREAD)
#define IMG_WIDTH 256
#define IMG_HEIGHT 256
#define IMG_SIZE (IMG_WIDTH*IMG_HEIGHT)
#define BLOCKS_X (IMG_WIDTH/(THREADS_X*PIXELS_PER_THREAD))
#define BLOCKS_Y (IMG_HEIGHT/THREADS_Y)

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void convert2DEdgeImageTo1DArray( unsigned char const * const image,
unsigned int* const list, int* const glob_index ) {

unsigned int const x = blockIdx.x  * THREADS_X*PIXELS_PER_THREAD + threadIdx.x;
unsigned int const y = blockIdx.y  * THREADS_Y + threadIdx.y;

volatile int qindex = -1;
volatile __shared__ int sh_qindex[THREADS_Y];
volatile __shared__ int sh_qstart[THREADS_Y];
sh_qindex[threadIdx.y] = -1;

// Start by making an array
volatile __shared__ unsigned int sh_queue[MAX_QUEUE_LENGTH];

// Fill the queue
for(int i=0; i<PIXELS_PER_THREAD; i++)
{
    int const xx = i*THREADS_X + x;

    // Read one image pixel from global memory
    unsigned char const pixel = image[y*IMG_WIDTH + xx];
    unsigned int  const queue_val = (y << 16) + xx;

    if(pixel)
    {
        do {
            qindex++;
            sh_qindex[threadIdx.y] = qindex;
            sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + qindex] = queue_val;
        } while (sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + qindex] != queue_val);
    }

    // Reload index from smem (last thread to write to smem will have updated it)
    qindex = sh_qindex[threadIdx.y];
}

// Let thread 0 reserve the space required in the global list
__syncthreads();
if(threadIdx.x == 0 && threadIdx.y == 0)
{
    // Find how many items are stored in each list
    int total_index = 0;
    #pragma unroll
    for(int i=0; i<THREADS_Y; i++)
    {
        sh_qstart[i] = total_index;
        total_index += (sh_qindex[i] + 1u);
    }

    // Calculate the offset in the global list
    unsigned int global_offset = atomicAdd(glob_index, total_index);
    #pragma unroll
    for(int i=0; i<THREADS_Y; i++)
    {
        sh_qstart[i] += global_offset;
    }

}
__syncthreads();

// Copy local queues to global queue
/*
for(int i=0; i<=qindex; i+=THREADS_X)
{
    if(i + threadIdx.x > qindex)
        break;

    unsigned int qvalue = sh_queue[threadIdx.y*THREADS_X*PIXELS_PER_THREAD + i + threadIdx.x];
    list[sh_qstart[threadIdx.y] + i + threadIdx.x] = qvalue;
}
*/
}

void call2DTo1DKernel(unsigned char const * const h_image)
{
    // Device side allocation
    unsigned char *d_image = NULL;
    unsigned int *d_list = NULL;
    int h_nlist=0, *d_nlist = NULL;
    cudaMalloc((void**)&d_image, sizeof(unsigned char)*IMG_SIZE);
    cudaMalloc((void**)&d_list, sizeof(unsigned int)*IMG_SIZE);
    cudaMalloc((void**)&d_nlist, sizeof(int));
    cudaCheckErrors("cudamalloc fail");

    // Time measurement initialization
    cudaEvent_t start, stop, startio, stopio;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventCreate(&startio);
    cudaEventCreate(&stopio);
    float et, etio;

    // Start timer w/ io
    cudaEventRecord(startio,0);
    cudaMemcpy(d_nlist, &h_nlist, sizeof(int), cudaMemcpyHostToDevice);
    // Copy image data to device
    cudaMemcpy((void*)d_image, (void*)h_image, sizeof(unsigned char)*IMG_SIZE,    cudaMemcpyHostToDevice);
    cudaCheckErrors("cudamemcpy 1");
    // Start timer
    cudaEventRecord(start,0);

    // Kernel call
    // Phase 1 : Convert 2D binary image to 1D pixel array
    dim3 dimBlock1(THREADS_X, THREADS_Y);
    dim3 dimGrid1(BLOCKS_X, BLOCKS_Y);
    convert2DEdgeImageTo1DArray<<<dimGrid1, dimBlock1>>>(d_image, d_list, d_nlist);
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel fail");
    // Stop timer
    cudaEventRecord(stop,0);
    cudaEventSynchronize(stop);

    // Stop timer w/ io
    cudaEventRecord(stopio,0);
    cudaEventSynchronize(stopio);

    // Time measurement
    cudaEventElapsedTime(&et,start,stop);
    cudaEventElapsedTime(&etio,startio,stopio);

    // Time measurement deinitialization
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    cudaEventDestroy(startio);
    cudaEventDestroy(stopio);

    // Get list size
    cudaMemcpy((void*)&h_nlist, (void*)d_nlist, sizeof(int), cudaMemcpyDeviceToHost);
    cudaCheckErrors("cudaMemcpy 2");
    // Report on console
    printf("%d pixels processed...\n", h_nlist);

    // Device side dealloc
    cudaFree(d_image);
//    cudaFree(d_space);
    cudaFree(d_list);
    cudaFree(d_nlist);
}

int main(){

  unsigned char *image;

  image = (unsigned char *)malloc(IMG_SIZE * sizeof(unsigned char));
  if (image == 0) {printf("malloc fail\n"); return 0;}

  for (int i =0 ; i<IMG_SIZE; i++)
    image[i] = rand()%2;

  call2DTo1DKernel(image);
  call2DTo1DKernel(image);
  call2DTo1DKernel(image);
  call2DTo1DKernel(image);
  call2DTo1DKernel(image);
  cudaCheckErrors("some error");
  return 0;
}

$ nvcc -arch=sm_20 -O3 -o t216 t216.cu
$ ./t216
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
$ ./t216
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
32617 pixels processed...
$
Community
  • 1
  • 1
Robert Crovella
  • 143,785
  • 11
  • 213
  • 257