c++ - Output of cuda program is not what was expected

Question

#include<cuda_runtime.h>
#include<stdio.h>
#include<cuda.h>
#include<stdlib.h>


__global__ void setVal(char **c){

c[(blockIdx.y * gridDim.x) + blockIdx.x] = "hello";

}


int main(){

char **gpu = NULL;
cudaMalloc((void**)&gpu, 6 * sizeof(char *));
int i;
/*
I cannot access second level directly
for( i =0 ; i < 6 ;i++){
    cudaMalloc((void**)&gpu[i], 10 * sizeof(char));
}*/


dim3 grid(3,2);
setVal<<<grid, 1>>>(gpu);
char *p = (char*)malloc(10 * sizeof(char));
char *x[6];

cudaMemcpy(x, gpu, 6*sizeof(char*), cudaMemcpyDeviceToHost);
for( i =0 ; i< 6; i++){
    cudaMemcpy(p, x[i], 10*sizeof(char), cudaMemcpyDeviceToHost);
    //put synchronize here if problem
    printf("%s
",p);

}


getchar();
return 0;
}

Based on all the suggestions, i revised my code to make my concept correct. But, the code is still not working :(. Any help will be appreciated

Answer 1

Try this -- I tested it on a GTX 285 under CUDA 3.2 -- so it's a bit more restrictive than the current version, but it works.

#include<stdio.h>
#include<string.h>

__global__ void setValues(char** word)
{
    volatile char* myWord = word[blockIdx.x];

    myWord[0] = 'H';
    myWord[1] = 'o';
    myWord[2] = 'l';
    myWord[3] = 'a';
    myWord[4] = '';
}

int main()
{
    const size_t bufferSize = 32;
    const int nObjects = 10;

    char*  h_x[nObjects];
    char** d_x = 0;

    cudaMalloc( (void**)(&d_x), nObjects * sizeof(char*) );

    for ( int i=0; i < nObjects; i++ )
    {
        h_x[i] = NULL;
        cudaMalloc( (void**)(&h_x[i]), bufferSize * sizeof(char) );
        printf("h_x[%d] = %lx
",i,(unsigned long)h_x[i]);
    }

    cudaMemcpy( d_x, h_x, nObjects*sizeof(char*), cudaMemcpyHostToDevice);
    printf("Copied h_x[] to d_x[]
");

    char msg[] = "Hello World!";
    cudaMemcpy( h_x[0], msg, 13*sizeof(char), cudaMemcpyHostToDevice );

    /*  Force Thread Synchronization  */
    cudaError err = cudaThreadSynchronize();

    /*  Check for and display Error  */
    if ( cudaSuccess != err )
    {
        fprintf( stderr, "Cuda error in file '%s' in line %i : %s.
",
                __FILE__, __LINE__, cudaGetErrorString( err) );
    }

    setValues<<<nObjects,1>>>(d_x);

    /*  Force Thread Synchronization  */
    err = cudaThreadSynchronize();

    /*  Check for and display Error  */
    if ( cudaSuccess != err )
    {
        fprintf( stderr, "Cuda error in file '%s' in line %i : %s.
",
                __FILE__, __LINE__, cudaGetErrorString( err) );
    }

    printf("Kernel Completed Successfully.  Woot.

");

    char p[bufferSize];

    printf("d_x = %lx
", (unsigned long)d_x );
    printf("h_x = %lx
", (unsigned long)h_x );

    cudaMemcpy( h_x, d_x, nObjects*sizeof(char*), cudaMemcpyDeviceToHost);

    printf("d_x = %lx
", (unsigned long)d_x );
    printf("h_x = %lx
", (unsigned long)h_x );

    for ( int i=0; i < nObjects; i++ )
    {
        cudaMemcpy( &p, h_x[i], bufferSize*sizeof(char), cudaMemcpyDeviceToHost);
        printf("%d p[] = %s
",i,p);
    }

    /*  Force Thread Synchronization  */
    err = cudaThreadSynchronize();

    /*  Check for and display Error  */
    if ( cudaSuccess != err )
    {
        fprintf( stderr, "Cuda error in file '%s' in line %i : %s.
",
                __FILE__, __LINE__, cudaGetErrorString( err) );
    }

    getchar();

    return 0;
}

As @Jon notes, you can't pass x (as you had declared) it to the GPU, because it's an address which lives on the CPU. In the code above, I create an array of char*'s and pass them to a char** which I also allocated on the GPU. Hope this helps!