cuda-gdb 3.2 gives up the ghost upon build with enum in __global__ function

for cuda-gdb version :

cuda-gdb --version

NVIDIA (R) CUDA Debugger

3.2 release

Portions Copyright (C) 2008-2010 NVIDIA Corporation

GNU gdb 6.6

Copyright (C) 2006 Free Software Foundation, Inc.

GDB is free software, covered by the GNU General Public License, and you are

welcome to change it and/or distribute copies of it under certain conditions.

Type "show copying" to see the conditions.

There is absolutely no warranty for GDB.  Type "show warranty" for details.

This GDB was configured as "x86_64-unknown-linux-gnu".

cuda-gdb give up the ghost (crashes) when enumerated types are used in global functions. Last I checked enums were pretty basic of a type, just after primitive types, just before structs and well before classes. Simply try the code with and without the enumerated type in the global bitreverse call and remove from

#include <stdio.h>

#include <stdlib.h>

#include <matrixlib/gpu_matrix.h>

//#include <matrixlib/matrix_common.h>

// Simple 8-bit bit reversal Compute test

#define USE_MDO

enum MatrixDataOrderEnum

{

	MAT_NONE,

	MAT_ROW_MAJOR,

	MAT_COL_MAJOR

};

struct goop_struct

{

    int i;

    int j;

    int width;

};

__device__ bool foo( Matrix m )

{

	int y = 1234;

	m.width = 34;

	m.height = 1234;

	y = 345;

	y = 4564;

}

#define N 256

#ifdef USE_MDO

__global__ void bitreverse(unsigned int *data, MatrixDataOrderEnum mdo )

#else

__global__ void bitreverse(unsigned int *data  )

#endif

{

	data_t gpu_data[4];

	Matrix gpu_mat;

	gpu_mat.width = 23;

	//	gpu_matrix_init( gpu_data, 1, 4, sizeof( data_t ), COLUMN_MAJOR );

	MATRIX_INIT( gpu_mat, gpu_data, 1, 4, sizeof( data_t ), false );

	foo( gpu_mat );

	if( gpu_mat.width == 4 )

	{

		gpu_mat.height = 2;

	}

	else

	{

		gpu_mat.height = 123;

	}

	goop_struct goop;

goop.i = 23;

if( goop.i == 23 )

}

    {

        goop.j = 1;

    }

    else

    {

        goop.width = 456;

    }

unsigned int *idata = data;

    unsigned int x = idata[threadIdx.x];

    x = ((0xf0f0f0f0 & x) >> 4) | ((0x0f0f0f0f & x) << 4);

    x = ((0xcccccccc & x) >> 2) | ((0x33333333 & x) << 2);

    x = ((0xaaaaaaaa & x) >> 1) | ((0x55555555 & x) << 1);

    idata[threadIdx.x] = x;

}

int main(void)

{

    unsigned int *d = NULL; int i;

    unsigned int idata[N], odata[N];

    for (i = 0; i < N; i++)

    idata[i] = (unsigned int)i;

    cudaMalloc((void**)&d, sizeof(int)*N);

    cudaMemcpy(d, idata, sizeof(int)*N,

    cudaMemcpyHostToDevice);

MatrixDataOrderEnum mdo = MAT_ROW_MAJOR;

	#ifdef USE_MDO

    bitreverse<<<1, N>>>(d , mdo );

	#else

    bitreverse<<<1, N>>>(d );

	#endif

    cudaMemcpy(odata, d, sizeof(int)*N,

    cudaMemcpyDeviceToHost);

    for (i = 0; i < N; i++)

    printf("%u -> %u\n", idata[i], odata[i]);

    cudaFree((void*)d);

    return 0;

}

You have no idea how long it took me to find this problem in my code. ESPECIALLY WHEN IT COMPILES AND RUNS.

I clicked the post instead of preview before explaining more of the shear awesomeness of this problem. The crash of gdb also breaks on cudaMalloc, cudaMallloc3d, and friends. Yuppers it doesn’t seem to break on the call of the global function with the enum. It’s as though cuda-gdb gets lost when calling into the cudda util lib and there is a global function with an enum type compiled in the kernel code. Seemingly two total unrelated parts of the code. I even tried it with examples where the global function is not even called, but simply compiled in.

I cannot fully covey what a absolute and total colossal waste of time trying to figure this out has been for me. Thanks NVidia your the best!

Here’s a before and after for you NVidia:

with

#define USE_MDO

bdavis5@teradrive:~/projects/NIH2009/branches/trunk$ cuda-gdb ./build/Linux-2.6.32-27-server/install/bin/simple_cuda_app 

NVIDIA (R) CUDA Debugger

3.2 release

Portions Copyright (C) 2008-2010 NVIDIA Corporation

GNU gdb 6.6

Copyright (C) 2006 Free Software Foundation, Inc.

GDB is free software, covered by the GNU General Public License, and you are

welcome to change it and/or distribute copies of it under certain conditions.

Type "show copying" to see the conditions.

There is absolutely no warranty for GDB.  Type "show warranty" for details.

This GDB was configured as "x86_64-unknown-linux-gnu"...

Using host libthread_db library "/lib/libthread_db.so.1".

(cuda-gdb) b simple_cuda_app.cu:100

Breakpoint 1 at 0x4016ed: file /home/bdavis5/projects/NIH2009/branches/trunk/source/Matlab/lib/dsaLib/gpu/slightly_complicated_cuda_app/src/simple_cuda_app.cu, line 100.

(cuda-gdb) run

Starting program: /home/bdavis5/projects/NIH2009/branches/trunk/build/Linux-2.6.32-27-server/install/bin/simple_cuda_app 

[Thread debugging using libthread_db enabled]

[New process 6378]

[New Thread 140325206525728 (LWP 6378)]

[Switching to Thread 140325206525728 (LWP 6378)]

Breakpoint 1, main ()

    at /home/bdavis5/projects/NIH2009/branches/trunk/source/Matlab/lib/dsaLib/gpu/slightly_complicated_cuda_app/src/simple_cuda_app.cu:100

100         cudaMalloc((void**)&d, sizeof(int)*N);

(cuda-gdb) n

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6376_415992117.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6376_415992117.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6376_112457376.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6376_112457376.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6376_112457376.o

BACKTRACE (3 frames):

cuda-gdb[0x46379e]

/lib/libc.so.6(+0x33af0)[0x7f8750de7af0]

/lib/libc.so.6(+0x37ef08)[0x7f8751132f08]

bdavis5@teradrive:~/projects/NIH2009/branches/trunk$

Above crashes cuda-gdb and I am returned to the prompt

bdavis5@teradrive:~/projects/NIH2009/branches/trunk$ cuda-gdb ./build/Linux-2.6.32-27-server/install/bin/simple_cuda_app 

NVIDIA (R) CUDA Debugger

3.2 release

Portions Copyright (C) 2008-2010 NVIDIA Corporation

GNU gdb 6.6

Copyright (C) 2006 Free Software Foundation, Inc.

GDB is free software, covered by the GNU General Public License, and you are

welcome to change it and/or distribute copies of it under certain conditions.

Type "show copying" to see the conditions.

There is absolutely no warranty for GDB.  Type "show warranty" for details.

This GDB was configured as "x86_64-unknown-linux-gnu"...

Using host libthread_db library "/lib/libthread_db.so.1".

(cuda-gdb) b simple_cuda_app.cu:99

Breakpoint 1 at 0x4016ad: file /home/bdavis5/projects/NIH2009/branches/trunk/source/Matlab/lib/dsaLib/gpu/slightly_complicated_cuda_app/src/simple_cuda_app.cu, line 99.

(cuda-gdb) r

Starting program: /home/bdavis5/projects/NIH2009/branches/trunk/build/Linux-2.6.32-27-server/install/bin/simple_cuda_app 

[Thread debugging using libthread_db enabled]

[New process 6672]

[New Thread 139770511484704 (LWP 6672)]

[Switching to Thread 139770511484704 (LWP 6672)]

Breakpoint 1, main ()

    at /home/bdavis5/projects/NIH2009/branches/trunk/source/Matlab/lib/dsaLib/gpu/slightly_complicated_cuda_app/src/simple_cuda_app.cu:99

99          cudaMalloc((void**)&d, sizeof(int)*N);

(cuda-gdb) s

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_219852114.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_2071432535.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_2071432535.o

warning: no loadable sections found in added symbol-file /tmp/tmp_cudagdb_elf_28503_6662_2071432535.o

100         cudaMemcpy(d, idata, sizeof(int)*N,

(cuda-gdb) s

103         MatrixDataOrderEnum mdo = MAT_ROW_MAJOR;

(cuda-gdb)

Hey this is successful. Who knew enum types could be so difficult for a debugger.

You are absolutely right enums in the device code is pretty basic stuff but somehow slipped through all the testing coverage. We did hear from one of the customer directly so this issue is now fixed and will be available in the 4.0 RC build end of Feb. In future you can reach the debugger team directly by emailing @ cuda-debugger-bugs@nvidia.com