# What's wrong with this? (malign-double?, passing integer array?, bug?)

**URL:** <https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401>\
**Category:** CUDA Programming and Performance\
**Created:** [August 21, 2010, 12:37am UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401 "2010-08-21T00:37:28Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![pcasto](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pcasto](https://forums.developer.nvidia.com/u/pcasto)\
**Post date:** [August 21, 2010, 12:37am UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401/1 "2010-08-21T00:37:28Z")

</div>

Hello everyone,

I was having trouble with one of my kernel in my CFD code. I created this simple code that illustrates the problem I have.

[codebox]

#include \<stdlib.h\>

#include \<stdio.h\>

#include

#include

#include

#include \<cuda.h\>

#include \<string.h\>

**global** void kernel1(float \*g\_array\_out, int \*g\_table, int dummy1, int dummy2 )

{

```
            // ints

            int j,index_l;

            int b_pid, g_pid;

            int tid = threadIdx.x;

```

//int k=0;

```
            for(int k =0; k < dummy1; k++)

            {

                    b_pid = k*blockDim.x + tid;

                    g_pid = blockIdx.x*dummy1*blockDim.x + b_pid;

```

if(g\_pid \< dummy2)

```
                    {

                            index_l = g_table[g_pid];

                            //index_l = g_pid;

```

for(j = 0; j \< 4; j++)

```
                            {

                                    g_array_out[index_l] = 1.;

                                    index_l += dummy2;

                            }

```

}

```
            }

```

}

int main()

{

```
    float *h_array_out= 0;

    float *d_array_out= 0;

```

int \*d\_table = 0;

```
    int *h_table = 0;

```

int block\_size = 16;

```
    int grid_size = 1;

```

// allocate look\_up table

```
    cudaMalloc((void**) &d_table, 16*sizeof(int));

    cudaMemset(d_table,0, 16*sizeof(int));

    h_table = (int*)malloc( 16*sizeof(int));

```

for(int g\_pid=0;g\_pid\<16;g\_pid++)

```
    {

            h_table[g_pid] = g_pid;

    }

```

cudaMemcpy(d\_table, h\_table, 16\*sizeof(int), cudaMemcpyHostToDevice);

```
    free(h_table);

```

// allocate array\_out

```
    cudaMalloc((void**) &d_array_out, sizeof(float)*4*16);

    cudaMemset(d_array_out,0,sizeof(float)*4*16);

```

if (d\_array\_out== 0 )

```
    {

            printf("couldn't allocate device memory\n");

            exit(1);

    }

```

kernel1\<\<\<grid\_size,block\_size\>\>\>(d\_array\_out,d\_table,1,16);

h\_array\_out= (float\*)malloc( 4_16_sizeof(float));

```
    cudaMemcpy(h_array_out,d_array_out,sizeof(float)*4*16,cudaMe

```

mcpyDeviceToHost);

for(int i=0;i\<4\*16;i++)

```
    {

            printf("%f\n",h_array_out[i]);

    }

    free(h_array_out);

```

cudaFree(d\_array\_out);

```
    cudaFree(d_table);

```

}

[/codebox]

I would expect the output array (array\_out) to be an array full of ones. However, when I compile with:

nvcc -arch=sm\_13 -I/usr/local/cuda/include …/src/cuda\_methods.cu -o cuda\_exec

and execute the program, the “array\_out” array has zeros in positions [16:47]. What is also really strange is that if I comment the outer loop in the kernel (since dummy1 is equal to 1 anyway), then I obtain what I expect: array\_out full of ones. Also, if I replace the line: “index\_l = d\_table[g\_pid]” by “index\_l = g\_pid”, I also get the good answer, even though d\_table[g\_pid] = g\_pid. It seems like it’s something related to integer arrays not being dereferenced properly. I heard about the “-malign-double” option but that doesn’t seem to work when used with nvcc.

Do you guys have any idea what’s going on?

I would really appreciate your help.

---

<div class="post-metadata">

**Author:** ![pcasto](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pcasto](https://forums.developer.nvidia.com/u/pcasto)\
**Post date:** [August 24, 2010, 7:11pm UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401/2 "2010-08-24T19:11:51Z")

</div>

Update: If I compile with the -m32 flag, then I get the results I expect. In the Cuda 3.1 release note, in the “Known Issues” section, I found the following:

When compiling with GCC, special care must be taken for structs that

contain 64-bit integers. This is because GCC aligns long longs

to a 4 byte boundary by default, while NVCC aligns long longs

to an 8 byte boundary by default. Thus, when using GCC to

compile a file that has a struct/union, users must give the

-malign-double

option to GCC. When using NVCC, this option is automatically

passed to GCC.

However, when I compile with the “–verbose” option, without the -m32 option, I don’t see the “malign-double” option being passed to gcc or g++. Here is the output

[codebox]

#$ _SPACE_=

#$ _CUDART_=cudart

#$ _HERE_=/usr/local/cuda/bin

#$ _THERE_=/usr/local/cuda/bin

#$ _TARGET\_SIZE_=64

#$ TOP=/usr/local/cuda/bin/…

#$ LD\_LIBRARY\_PATH=/usr/local/cuda/bin/…/lib:/usr/local/cuda/bin/…/extools/lib:/usr/local/cuda/lib:/usr/local/cuda/lib64:/usr/lib64:/usr/lib64/mpich2/lib:/usr/local/cuda/lib:/usr/local/cuda/lib64:

#$ PATH=/usr/local/cuda/bin/…/open64/bin:/usr/local/cuda/bin:/usr/local/tecplot/bin:/usr/local/cuda/bin:/usr/lib64/mpich2/bin:/home/pcasto/mpich2-install/bin:/usr/local/cuda/bin:/usr/lib64/qt-3.3/bin:/usr/kerberos/sbin:/usr/kerberos/bin:/usr/lib64/ccache:/usr/local/bin:/usr/bin:/bin:/usr/local/sbin:/usr/sbin:/sbin:/home/pcasto/bin

#$ INCLUDES=“-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart”

#$ LIBRARIES= “-L/usr/local/cuda/bin/…/lib64” -lcudart

#$ CUDAFE\_FLAGS=

#$ OPENCC\_FLAGS=

#$ PTXAS\_FLAGS=

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c++ -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -include “cuda\_runtime.h” -m64 -o “/tmp/tmpxft\_00000662\_00000000-4\_cuda\_methods.cpp1.ii” “…/src/cuda\_methods.cu”

#$ cudafe --m64 --gnu\_version=40404 -tused --no\_remove\_unneeded\_entities --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.c” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.stub.c” --gen\_device\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.gpu” --include\_file\_name “/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “/tmp/tmpxft\_00000662\_00000000-4\_cuda\_methods.cpp1.ii”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-5\_cuda\_methods.cpp2.i” “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.gpu”

#$ cudafe --m64 --gnu\_version=40404 --c --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.c” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.stub.c” --gen\_device\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.gpu” --include\_file\_name “/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “/tmp/tmpxft\_00000662\_00000000-5\_cuda\_methods.cpp2.i”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDABE\_\_ -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.gpu”

#$ filehash -s " " “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” \> “/tmp/tmpxft\_00000662\_00000000-8\_cuda\_methods.hash”

#$ nvopencc -TARG:compute\_13 -m64 -CG:ftz=1 -CG:prec\_div=0 -CG:prec\_sqrt=0 “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” -o “/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx”

#$ ptxas -arch=sm\_13 -m64 “/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx” -o “/tmp/tmpxft\_00000662\_00000000-9\_cuda\_methods.sm\_13.cubin”

#$ fatbin --key=“ffed6152a4eff03f” --source-name=“…/src/cuda\_methods.cu” --usage-mode=" " --embedded-fatbin=“/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “–image=profile=compute\_13,file=/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx” “–image=profile=sm\_13,file=/tmp/tmpxft\_00000662\_00000000-9\_cuda\_methods.sm\_13.cubin”

#$ gcc -E -x c++ “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -include “cuda\_runtime.h” -m64 -o “/tmp/tmpxft\_00000662\_00000000-10\_cuda\_methods.cpp4.ii” “…/src/cuda\_methods.cu”

#$ cudafe++ --m64 --gnu\_version=40404 --parse\_templates --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.cpp” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.stub.c” “/tmp/tmpxft\_00000662\_00000000-10\_cuda\_methods.cpp4.ii”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c++ -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-11\_cuda\_methods.ii” “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.cpp”

#$ gcc -c -x c++ “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -fpreprocessed -m64 -o “/tmp/tmpxft\_00000662\_00000000-12\_cuda\_methods.o” “/tmp/tmpxft\_00000662\_00000000-11\_cuda\_methods.ii”

#$ g++ -m64 -o “cuda\_methods” -Wl,–start-group “/tmp/tmpxft\_00000662\_00000000-12\_cuda\_methods.o” “-L/usr/local/cuda/bin/…/lib64” -lcudart -Wl,–end-group

[/codebox]

The code I posted is quite simple, yet I can’t get it to compile and give me the expected result. Has anyone tried to compile and run the code? I would be curious to see if you’re getting the same result. I would really appreciate any help.

---

<div class="post-metadata">

**Author:** ![pcasto](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pcasto](https://forums.developer.nvidia.com/u/pcasto)\
**Post date:** [August 24, 2010, 7:11pm UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401/3 "2010-08-24T19:11:51Z")

</div>

Update: If I compile with the -m32 flag, then I get the results I expect. In the Cuda 3.1 release note, in the “Known Issues” section, I found the following:

When compiling with GCC, special care must be taken for structs that

contain 64-bit integers. This is because GCC aligns long longs

to a 4 byte boundary by default, while NVCC aligns long longs

to an 8 byte boundary by default. Thus, when using GCC to

compile a file that has a struct/union, users must give the

-malign-double

option to GCC. When using NVCC, this option is automatically

passed to GCC.

However, when I compile with the “–verbose” option, without the -m32 option, I don’t see the “malign-double” option being passed to gcc or g++. Here is the output

[codebox]

#$ _SPACE_=

#$ _CUDART_=cudart

#$ _HERE_=/usr/local/cuda/bin

#$ _THERE_=/usr/local/cuda/bin

#$ _TARGET\_SIZE_=64

#$ TOP=/usr/local/cuda/bin/…

#$ LD\_LIBRARY\_PATH=/usr/local/cuda/bin/…/lib:/usr/local/cuda/bin/…/extools/lib:/usr/local/cuda/lib:/usr/local/cuda/lib64:/usr/lib64:/usr/lib64/mpich2/lib:/usr/local/cuda/lib:/usr/local/cuda/lib64:

#$ PATH=/usr/local/cuda/bin/…/open64/bin:/usr/local/cuda/bin:/usr/local/tecplot/bin:/usr/local/cuda/bin:/usr/lib64/mpich2/bin:/home/pcasto/mpich2-install/bin:/usr/local/cuda/bin:/usr/lib64/qt-3.3/bin:/usr/kerberos/sbin:/usr/kerberos/bin:/usr/lib64/ccache:/usr/local/bin:/usr/bin:/bin:/usr/local/sbin:/usr/sbin:/sbin:/home/pcasto/bin

#$ INCLUDES=“-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart”

#$ LIBRARIES= “-L/usr/local/cuda/bin/…/lib64” -lcudart

#$ CUDAFE\_FLAGS=

#$ OPENCC\_FLAGS=

#$ PTXAS\_FLAGS=

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c++ -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -include “cuda\_runtime.h” -m64 -o “/tmp/tmpxft\_00000662\_00000000-4\_cuda\_methods.cpp1.ii” “…/src/cuda\_methods.cu”

#$ cudafe --m64 --gnu\_version=40404 -tused --no\_remove\_unneeded\_entities --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.c” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.stub.c” --gen\_device\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.gpu” --include\_file\_name “/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “/tmp/tmpxft\_00000662\_00000000-4\_cuda\_methods.cpp1.ii”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-5\_cuda\_methods.cpp2.i” “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.gpu”

#$ cudafe --m64 --gnu\_version=40404 --c --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.c” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.stub.c” --gen\_device\_file\_name “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.gpu” --include\_file\_name “/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “/tmp/tmpxft\_00000662\_00000000-5\_cuda\_methods.cpp2.i”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDABE\_\_ -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” “/tmp/tmpxft\_00000662\_00000000-6\_cuda\_methods.cudafe2.gpu”

#$ filehash -s " " “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” \> “/tmp/tmpxft\_00000662\_00000000-8\_cuda\_methods.hash”

#$ nvopencc -TARG:compute\_13 -m64 -CG:ftz=1 -CG:prec\_div=0 -CG:prec\_sqrt=0 “/tmp/tmpxft\_00000662\_00000000-7\_cuda\_methods.cpp3.i” -o “/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx”

#$ ptxas -arch=sm\_13 -m64 “/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx” -o “/tmp/tmpxft\_00000662\_00000000-9\_cuda\_methods.sm\_13.cubin”

#$ fatbin --key=“ffed6152a4eff03f” --source-name=“…/src/cuda\_methods.cu” --usage-mode=" " --embedded-fatbin=“/tmp/tmpxft\_00000662\_00000000-3\_cuda\_methods.fatbin.c” “–image=profile=compute\_13,file=/tmp/tmpxft\_00000662\_00000000-2\_cuda\_methods.ptx” “–image=profile=sm\_13,file=/tmp/tmpxft\_00000662\_00000000-9\_cuda\_methods.sm\_13.cubin”

#$ gcc -E -x c++ “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDACC\_\_ -C -include “cuda\_runtime.h” -m64 -o “/tmp/tmpxft\_00000662\_00000000-10\_cuda\_methods.cpp4.ii” “…/src/cuda\_methods.cu”

#$ cudafe++ --m64 --gnu\_version=40404 --parse\_templates --gen\_c\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.cpp” --stub\_file\_name “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.stub.c” “/tmp/tmpxft\_00000662\_00000000-10\_cuda\_methods.cpp4.ii”

#$ gcc -D\_\_CUDA\_ARCH\_\_=130 -E -x c++ -DCUDA\_DOUBLE\_MATH\_FUNCTIONS “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -D\_\_CUDA\_FTZ -m64 -o “/tmp/tmpxft\_00000662\_00000000-11\_cuda\_methods.ii” “/tmp/tmpxft\_00000662\_00000000-1\_cuda\_methods.cudafe1.cpp”

#$ gcc -c -x c++ “-I/usr/local/cuda/bin/…/include” “-I/usr/local/cuda/bin/…/include/cudart” -I. -fpreprocessed -m64 -o “/tmp/tmpxft\_00000662\_00000000-12\_cuda\_methods.o” “/tmp/tmpxft\_00000662\_00000000-11\_cuda\_methods.ii”

#$ g++ -m64 -o “cuda\_methods” -Wl,–start-group “/tmp/tmpxft\_00000662\_00000000-12\_cuda\_methods.o” “-L/usr/local/cuda/bin/…/lib64” -lcudart -Wl,–end-group

[/codebox]

The code I posted is quite simple, yet I can’t get it to compile and give me the expected result. Has anyone tried to compile and run the code? I would be curious to see if you’re getting the same result. I would really appreciate any help.

---

<div class="post-metadata">

**Author:** ![Cliff\_Woolley](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/cliff_woolley/32/14043_2.png) [@Cliff\_Woolley](https://forums.developer.nvidia.com/u/Cliff_Woolley)\
**Post date:** [August 28, 2010, 6:05pm UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401/4 "2010-08-28T18:05:22Z")

</div>

Hi pcasto,

Thanks for your report. I’ll try to reproduce this issue and report a bug if it’s not already fixed.

–Cliff

---

<div class="post-metadata">

**Author:** ![Cliff\_Woolley](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/cliff_woolley/32/14043_2.png) [@Cliff\_Woolley](https://forums.developer.nvidia.com/u/Cliff_Woolley)\
**Post date:** [August 28, 2010, 6:05pm UTC](https://forums.developer.nvidia.com/t/whats-wrong-with-this-malign-double-passing-integer-array-bug/18401/5 "2010-08-28T18:05:22Z")

</div>

Hi pcasto,

Thanks for your report. I’ll try to reproduce this issue and report a bug if it’s not already fixed.

–Cliff