Debugging CUDA - CudaUnknownError

Question

I am trying to use CUDA to create a bitmap image of the mandlebrot set. I have looked at a few tutorials and already got some help here for the process of integrating the unmanaged CUDA dll with the managed C# gui. The problem I am having now is that my CUDA dll is not forming the bitmap correctly - and when I use an error checking macro on cudaDeviceSynchronize() after the kernel launch, I get cudaUnknownError.

Here is the relevant code:

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}



struct complex
{
    float r, i;

    __device__ complex(float _r, float _i) : r(_r), i(_i) {}

    __device__ float magnitudeSquared(){ return (r*r + i*i) ; }

    __device__ complex& operator*=(const complex& rhs)
    {
        r = (r * rhs.r - i * rhs.i);
        i = (r * rhs.i + i * rhs.r);
        return *this;
    }

    __device__ complex& operator+=(const complex& rhs)
    {
        r = (r + rhs.r);
        i = (i + rhs.i);
        return *this;
    }
};

__device__ int mandlebrotDiverge(complex *z)
{
    complex c(*z);
    int i = 0;
    for(i = 0; i < MAX_ITERATIONS; i++)
    {
        *z *= *z;
        *z += c;
        if(z->magnitudeSquared() >= 2)
        {
            return 1;
        }
    }
    return 0;

}

__global__ void kernel(int *ptr, int width, int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    float scale = 1.5f;
    complex z(scale*(float)(width/2 - x)/(width/2), scale*(float)(height/2 - y)/(height/2));

    if(offset < (1920*1080))
    {
        int mValue = mandlebrotDiverge(&z);
        ptr[offset*3 + (uint8_t)0] = (uint8_t)(mValue*255);
        ptr[offset*3 + (uint8_t)1] = (uint8_t)(mValue*255);
        ptr[offset*3 + (uint8_t)2] = (uint8_t)(mValue*255);
    }
}




extern "C" __declspec(dllexport) void __cdecl generateBitmap(void *bitmap)
{
    int width = 1920;
    int height = 1080;
    int *dev_bmp;

    dim3 blocks(width/16, height/16);
    dim3 threads(16, 16);

    gpuErrchk(cudaMalloc((void**)&dev_bmp, (3*width*height)));

    kernel<<<blocks, threads>>>(dev_bmp, width, height);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(bitmap, dev_bmp, (width*height*3), cudaMemcpyDeviceToHost));
    cudaFree(dev_bmp);
}

When I step through the code, everything appears to be working correctly until I get to gpuErrchk(cudaDeviceSynchronize()); - when I step into that, the error code simply says 'cudaUnknownError'. I have no clue what I am doing wrong at this point. Any help or advice on how to improve this solution would be appreciated.

EDIT:

Okay, looked at CUDA_memcheck, and I get this error (hundreds of times):

========= CUDA-MEMCHECK
========= Invalid __global__ write of size 4
=========     at 0x00000a10 in C:/.../kernel.cu:77:kernel(int*, int, int)
=========     by thread (15,11,0) in block (1,17,0)
=========     Address 0x05a37f74 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time

So I changed from *int to *unsigned char because I am trying to allocate arrays of individual bytes, not ints. Lots of errors cleared up, but now I get this:

========= CUDA-MEMCHECK
========= Program hit error 6 on CUDA API call to cudaDeviceSynchronize 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:C:\Windows\system32\nvcuda.dll (cuD3D11CtxCreate + 0x102459) [0x11e4b9]
=========     Host Frame:C:\...\cudart32_55.dll (cudaDeviceSynchronize + 0xdd) [0x1149d]
=========     Host Frame:C:\...\FractalMaxUnmanaged.dll (generateBitmap + 0xf0) [0x97c0]
=========
========= ERROR SUMMARY: 1 error

Okay, I'm making progress, but now when I step through the c# application, the byte buffer has 255 as the value for every byte, which doesn't make sense. Here is the c# code:

public unsafe class NativeMethods
{

    [DllImport(@"C:\Users\Bill\Documents\Visual Studio 2012\Projects\FractalMaxUnmanaged\Debug\FractalMaxUnmanaged.dll", CallingConvention=CallingConvention.Cdecl)]
    public static extern void generateBitmap(void *bitmap);

    public static Bitmap create()
    {
        byte[] buf = new byte[1920 * 1080 * 3];
        fixed (void* pBuffer = buf)
        {
            generateBitmap(pBuffer);

        }
        IntPtr unmanagedPtr = Marshal.AllocHGlobal(buf.Length);
        Marshal.Copy(buf, 0, unmanagedPtr, buf.Length);
        Bitmap img = new Bitmap(1920, 1080, 1920*3, PixelFormat.Format24bppRgb, unmanagedPtr);

        Marshal.FreeHGlobal(unmanagedPtr);

        return img;
    }
}

Answer 1

Your problem here is that your memory allocations and copies are wrong, you are forgetting that cudaMalloc/cudaMemcpy expect the size in bytes. Since int uses 4 bytes, you are actually allocating less memory than required by your kernel. Use this instead (or use unsigned char which only needs 1 byte):

cudaMalloc((void**)&dev_bmp, (3*width*height)*sizeof(int));
cudaMemcpy(bitmap, dev_bmp, (3*width*height)*sizeof(int), cudaMemcpyDeviceToHost);

Also make sure that bitmap was allocated correctly. As @Eugene said, using cuda-memcheck is a good way to find the source of that kind of errors.