Problem occurred when compiling julia_gpu.cu from Cuda by example using nvcc

Accelerated Computing CUDA CUDA Programming and Performance
#1

A problem occurred when I was compiling julia_gpu.cu from Cuda by example using nvcc. I guess the problem was caused by compiler. Dose anybody know how to solve it ?
PS: I attached the codes and error info below:

julia_gpu.cu :

#include "../common/book.h"
#include "../common/cpu_bitmap.h"

#define DIM 1000

struct cuComplex {
    float   r;
    float   i;
    cuComplex( float a, float b ) : r(a), i(b)  {}
    __device__ float magnitude2( void ) {
        return r * r + i * i;
    }
    __device__ cuComplex operator*(const cuComplex& a) {
        return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
    }
    __device__ cuComplex operator+(const cuComplex& a) {
        return cuComplex(r+a.r, i+a.i);
    }
};

__device__ int julia( int x, int y ) {
    const float scale = 1.5;
    float jx = scale * (float)(DIM/2 - x)/(DIM/2);
    float jy = scale * (float)(DIM/2 - y)/(DIM/2);

    cuComplex c(-0.8, 0.156);
    cuComplex a(jx, jy);

    int i = 0;
    for (i=0; i<200; i++) {
        a = a * a + c;
        if (a.magnitude2() > 1000)
            return 0;
    }

    return 1;
}

__global__ void kernel( unsigned char *ptr ) {
    // map from blockIdx to pixel position
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;

    // now calculate the value at that position
    int juliaValue = julia( x, y );
    ptr[offset*4 + 0] = 255 * juliaValue;
    ptr[offset*4 + 1] = 0;
    ptr[offset*4 + 2] = 0;
    ptr[offset*4 + 3] = 255;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
};

int main( void ) {
    DataBlock   data;
    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char    *dev_bitmap;

    HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap, bitmap.image_size() ) );
    data.dev_bitmap = dev_bitmap;

    dim3    grid(DIM,DIM);
    kernel<<<grid,1>>>( dev_bitmap );

    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );
                              
    HANDLE_ERROR( cudaFree( dev_bitmap ) );
                              
    bitmap.display_and_exit();
}

…/common/cpu_bitmap.h :

//

#ifndef __CPU_BITMAP_H__
#define __CPU_BITMAP_H__

#include "gl_helper.h"

struct CPUBitmap {
    unsigned char    *pixels;
    int     x, y;
    void    *dataBlock;
    void (*bitmapExit)(void*);

    CPUBitmap( int width, int height, void *d = NULL ) {
        pixels = new unsigned char[width * height * 4];
        x = width;
        y = height;
        dataBlock = d;
    }

    ~CPUBitmap() {
        delete [] pixels;
    }

    unsigned char* get_ptr( void ) const   { return pixels; }
    long image_size( void ) const { return x * y * 4; }

    void display_and_exit( void(*e)(void*) = NULL ) {
        CPUBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        bitmapExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c=1;
        char* dummy = "";
        glutInit( &c, &dummy );
        glutInitDisplayMode( GLUT_SINGLE | GLUT_RGBA );
        glutInitWindowSize( x, y );
        glutCreateWindow( "bitmap" );
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        glutMainLoop();
    }

     // static method used for glut callbacks
    static CPUBitmap** get_bitmap_ptr( void ) {
        static CPUBitmap   *gBitmap;
        return &gBitmap;
    }

   // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
            case 27:
                CPUBitmap*   bitmap = *(get_bitmap_ptr());
                if (bitmap->dataBlock != NULL && bitmap->bitmapExit != NULL)
                    bitmap->bitmapExit( bitmap->dataBlock );
                exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw( void ) {
        CPUBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor( 0.0, 0.0, 0.0, 1.0 );
        glClear( GL_COLOR_BUFFER_BIT );
        glDrawPixels( bitmap->x, bitmap->y, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
        glFlush();
    }
};

#endif  // __CPU_BITMAP_H__

Error Info:

../common/cpu_bitmap.h(49): warning: conversion from a string literal to "char *" is deprecated
#2

That’s a warning, not an error.

If you google this warning message, you’ll find others who have run into it, it has nothing to do with CUDA. What’s going on is the language (C++) is evolving, and the newer compilers are evolving with it.

I believe you should be able to fix it by changing this:

char* dummy = "";

to this:

const char* dummy = "";
#3

No, it does not work. Because the problem occurred in a .h file(cpu_bitmap.h), which is difficult to be modified.
I think the problem is nvcc assigned a wrong compiler for the .h file(should be gcc, but nvcc used g++).
Do you know how to solve it ?

PS: The codes above are copied from CUDA by Example , which is an introductory tutorials from NVIDIA. So I think they’re unlikely wrong. The problem is probably in the compiling process(missing some parameters when compiling the codes with nvcc ?)

#4

The CUDA language is a derivative of C++, so it is entirely appropriate that the host code is passed to g++.

In recent years GNU compilers have added various warnings for things they did not use to warn about. In addition, although my recollection is hazy, I am fairly certain that at some point in the past they changed string literals from type char * to const char *. The book “CUDA by Example” came out five years ago, and obviously it was not possible for its authors to anticipate what kind of warnings GNU compilers may emit in the future.

You can either ignore the warning (nothing bad will happen), add a compiler flag to your build that suppresses this particular warning, or change the header file to add the const qualifier to the variable definition in question. I am afraid I don’t understand why you say that the header file is difficult to modify.

#5

I completely agree with you. The book “CUDA by Example” came out five years ago, and that’s why the newest CUDA compiler reported some warnings which may not be reported by a compiler developed 5 years ago.

Unfortunately, however, the char * variable defined by the header file cpu_bitmap.h is used for many times and in many places in the header file. Actually, I have tried to change char * to const char * but a lot more warnings and errors appeared. I think I can handle them by looking deeply in the codes and spending time in carefully modifying a lot of lines in the header file. But this kind of work would waste much time which is so precious for me. So I want to find a quick-and-convenient solution and it will help me a lot.

OK, I admit that I am just too lazy. Please stop laughing at me :)

#6

If use ‘const char* dummy = “”;’, an error will be reported like this:

../common/cpu_bitmap.h(50): error: argument of type "const char **" is incompatible with parameter of type "char **"
#7
#8

That error makes sense based on the definition of const. I am well familiar with the frustration that results from trying to address warnings that a new compiler version adds by default without the programmer having asked for it. Trying to apply a quick fix then often leads to warnings or errors elsewhere. Best I can tell without diving into the details it is actually the host compiler giving the warning, not nvcc itself.

My recommendation would be: Since your goal is presumably learning about CUDA, ignore the warning and focus on the CUDA-specific elements of the code. Presumably the example code never writes through the pointer dummy, thereby violating the read-only properties of the string literal. So nothing bad will happen from ignoring the warning in this case.

[Later:]

What happens if you make a change like so:

char* dummy = (char *)(void *)"";

Basically this “casts away const-ness”, and the intermediate (void *) cast may cause g++ to ignore this.

1 Like
#9

Unfortunately, however, the bad thing is I can’t get an executable file from a lot of codes who used that header file. Therefore I cannot check whether those codes are correct or learn from modifying and running those codes. That’s very annoying for a newbie in CUDA programming.

#10

Warnings only prevent an executable from being generated if the build uses a compiler switch that instructs the compiler to treat warnings as errors (for gcc/g++ that switch is -Werror, I think). You do not have to build with that compiler flag, although in general that is good practice especially in larger projects.

I updated my previous reply with a potential workaround while you posted message #9. That’s the best idea I have right now for a local fix.

#11

After I made the change you provided, the warnings were gone. And I checked the codes again and found that even though I don’t make the change, the codes can still be compiled and linked to an executable file except that warnings appear. It turned out that the problem that no executable was output was caused by another error…

So you are right! Thanks a lot for sharing your knowledge and skill with me.

And can you explain how your change to the codes works ?

#12

The type qualifier const in C/C++ is a data attribute that means “read only”. A pointer to a read-only char looks like so:

const char *cp1; // usual idiom: "cp1 is pointer to a `char` that is read-only"
char const *cp2; // alternative idiom: "cp2 is pointer to a read-only `char`"

In the case at hand, literal string constants are placed in read-only storage by the compiler, and the type of “” is therefore const char *. The compiler complains when a pointer to a read-only object is assigned to a pointer to writable storage, as that pointer could now be used to write to the read-only object, which is obviously not something we want to happen:

const char *cp1 = "string";  // read-only data object
char * cp2 = cp1;            // implicit cast
char * cp3 = (char *)cp1;    // explicit cast
*cp2 =  // oh no, write access to read-only object!
*cp3 =  // oh no, write access to read-only object!

Using C-programmer jargon, the explicit cast that generates cp3 is said to “cast away constness”.

The idiom of using an intermediate cast to a generic pointer, that is void *, convinces the compiler that the programmer knows what they are doing, thus suppressing the warning. This may not work with every compiler, as compilers have substantial freedom when it comes to emitting warnings; It works with all GNU compilers though as far as I recall.

The same technique works with data that has alignment restrictions, when you cast a pointer to a type with less strict alignment requirement to a pointer to a type with more stringent alignment requirements, which can break the code if the hardware requires all accesses to be aligned to the width of the data type. NVIDIA GPUs for example have such a natural alignment requirement. The gcc/g++ flag -Wcast-align turns on warning about this issue. So when you use that flag and program

double *dp1;
double2 *dp2;
dp2 = (double2 *)dp1;
*dp1 =    // requires dp1 aligned to 8-byte boundary (address divisible by 8)
*dp2 =    // requires dp2 aligned to 16-byte boundary (address divisible by 16)

the compiler will warn about the assignment to dp2. To avoid the warning when you can guarantee, by construction, that dp2 will always be suitably aligned, you would use this instead:

dp2 = (double2 *)(void *)dp1;
#13

[/quote]

I really learned a lot from your posts, and thanks again!

I am a bioinformatician and I don’t have very strong background in computer science. But I have strong desire to use the cutting-edge computing techniques to finish my work. You know, the development of DNA sequencing technology has outpaced Moore-law by 3-4 times. So a huge amount of sequence data is approaching and I have to use more powerful technique, say GPU, to analyze data, to build models and to write efficient programs.

So maybe I will become a good CUDA programmer in future :)

#14

I am glad to hear my explanations have been somewhat helpful. I am a computer science guy by degree, and know very little about bioinformatics :-) I have been following the application of GPU computing to bio informatics somewhat, and at one time seriously considered hiring on with a company that makes fast DNA sequencers. Fascinating science and technology, with great strides being made all the time. As far as I can tell, GPU computing and CUDA have proven valuable in successfully tackling various sub-problems in the overall “sequencing universe”.

My experience over the years in working with customers adopting CUDA has generally been that it is much easier for a domain specialist to pick up CUDA knowledge than for a “CUDA guy” to pick up domain specific knowledge. That could be just me, but I would like to think that it has something to do with the fact that NVIDIA has put much effort into making the technology increasingly accessible and user friendly and has invested into education as well. Success and best of luck with your CUDA endeavours.