Wierd crash in modified VectorAdd example.

I have been experimenting with passing kernel arguments as global variables by initializing the variable using cudaMemcpyToSymbol. It worked at first. However, if in the class Arr, if i Just add an dummy data member after the float* myptr member, it crashes. What is even more perplexing is that moving the dummy data member below the Myptr data member resolves the problem.

I am on windows 7 64 bit , but the code is set to compile a 32 bit executable. I use a gtx 470 with driver 270.61, and cuda toolkit 3.2.

The problem file is attached below.

Debdatta Basu.
Tester.cu (4.23 KB)

Can’t view the file.

I have no experience with 64-bit system, but I’m pretty sure it’s a problem with your pointer size. You can download the 64-bit 4.0 toolkit and compile it in 64-bit. Hopefully that will solve the problems.

I have attached it below:


 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.


 * Please refer to the NVIDIA end user license agreement (EULA) associated

 * with this source code for terms and conditions that govern your use of

 * this software. Any use, reproduction, disclosure, or distribution of

 * this software and related documentation outside the terms of the EULA

 * is strictly prohibited.



/* Vector addition: C = A + B.


 * This sample is a very basic sample that implements element by element

 * vector addition. It is the same as the sample illustrating Chapter 3

 * of the programming guide with some additions like error checking.



// Includes

#include <stdio.h>

#include <cuda_runtime_api.h>

#include <cutil_inline.h>

// Variables

float* h_A;

float* h_B;

float* h_C;

float* d_A;

float* d_B;

float* d_C;

bool noprompt = false;

// Functions

void Cleanup(void);

void RandomInit(float*, int);

void ParseArguments(int, char**);

// Device code

class Arr



	__host__ void init(float* p)


		cudaMemcpyToSymbol(MyPtr, (void*)&p, sizeof(float*));


	//The bug is here. Remove the int* Dummy, or place it after the Myptr to get a correct result.

	int* Dummy;

	float* MyPtr;



__device__ Arr A;

__device__ Arr B;

__device__ Arr C;

__global__ void VecAdd(int N)



	int i = blockDim.x * blockIdx.x + threadIdx.x;

	if (i < N)

		C.MyPtr[i] = A.MyPtr[i] + B.MyPtr[i];


// Host code

int main(int argc, char** argv)


    printf("Vector addition\n");

    int N = 50000;

    size_t size = N * sizeof(float);

    ParseArguments(argc, argv);

// Allocate input vectors h_A and h_B in host memory

    h_A = (float*)malloc(size);

    if (h_A == 0) Cleanup();

    h_B = (float*)malloc(size);

    if (h_B == 0) Cleanup();

    h_C = (float*)malloc(size);

    if (h_C == 0) Cleanup();

// Initialize input vectors

    RandomInit(h_A, N);

    RandomInit(h_B, N);

// Allocate vectors in device memory

    cutilSafeCall( cudaMalloc((void**)&d_A, size) );

    cutilSafeCall( cudaMalloc((void**)&d_B, size) );

    cutilSafeCall( cudaMalloc((void**)&d_C, size) );

// Copy vectors from host memory to device memory

    cutilSafeCall( cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice) );

    cutilSafeCall( cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice) );

	//cudaMemcpyToSymbol(A.MyPtr, (void*)&d_A, sizeof(float*));

	//cudaMemcpyToSymbol(B.MyPtr, (void*)&d_B, sizeof(float*));

	//cudaMemcpyToSymbol(C.MyPtr, (void*)&d_C, sizeof(float*));




// Invoke kernel

    int threadsPerBlock = 256;

    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

VecAdd<<<blocksPerGrid, threadsPerBlock>>>(N);

    cutilCheckMsg("kernel launch failure");

#ifdef _DEBUG

    cutilSafeCall( cudaThreadSynchronize() );


// Copy result from device memory to host memory

    // h_C contains the result in host memory

    cutilSafeCall( cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost) );

// Verify result

    int i;

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

        float sum = h_A[i] + h_B[i];

        if (fabs(h_C[i] - sum) > 1e-5)



    printf("%s \n", (i == N) ? "PASSED" : "FAILED");



void Cleanup(void)


    // Free device memory

    if (d_A)


    if (d_B)


    if (d_C)


// Free host memory

    if (h_A)


    if (h_B)


    if (h_C)


cutilSafeCall( cudaThreadExit() );

if (!noprompt) {

        printf("\nPress ENTER to exit...\n");

        fflush( stdout);

        fflush( stderr);





// Allocates an array with random float entries.

void RandomInit(float* data, int n)


    for (int i = 0; i < n; ++i)

        data[i] = rand() / (float)RAND_MAX;


// Parse program arguments

void ParseArguments(int argc, char** argv)


    for (int i = 0; i < argc; ++i)

        if (strcmp(argv[i], "--noprompt") == 0 ||

			strcmp(argv[i], "-noprompt") == 0) 


            noprompt = true;




The pointer copying code doesn’t used hard coded pointer sizes. So that shouldn’t be a problem. I think this might be an alignment issue though Im not sure.

I got it to work! Simply change the Arr class in the above code to the following:

class Arr



	__host__ void init(float* p)


		int offset = 0;

		ALIGN_UP(offset, __alignof(int*));

		offset+= sizeof(int*);

		ALIGN_UP(offset, __alignof(float*));

		cudaMemcpyToSymbol(Dummy, &p, sizeof(float*), offset);


	//The bug is removed.


	int* Dummy;

	float* MyPtr;

} ;

All I have done is reference the Myptr in a roundabout way by specifying an offset from Dummy. It appears that the cudaMemcpyToSymbol call is having problems obtaining addresses to data members of the class other than the first member. Is this a bug?

-Debdatta Basu.


Sorry for that, but did anyone try it out? I want to get to the bottom of this issue.

Debdatta Basu.

How do you use MyPtr directly in host code when it’s supposed to exist only in device memory? Are you sure you are getting correct result?
I suppose in the second case you got it to run only because it’s no longer writing to address 0x00000000, but 0x00000004.

Either you do not put them as symbols or you get their symbol addresses before setting values in them.

I am getting the correct result. The manual says that cudamemcpytosymbol can take symbol names. Hence, I can use MyPtr directly.

const char*

that’s what it means by symbol name.