Am I using constant memory correctly?

Question

I wrote a simple Cuda C program which takes N numbers and multiplies them by a factor c; since this factor is a constant I decided to put it in the constant memory.

Here is the main.cu code:

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

__constant__ float c; // IS THIS OK?

typedef struct{
    float a;
}variable;
variable *var_CPU,*var_GPU;

#include "kernel.cu"

//==============================================================
int main(void){

   int i,N=100;
   var_CPU = (variable*)malloc(N*sizeof(variable));
   for(i=0;i<N;i++){
      var_CPU[i].a=(float)i;
      printf("var_CPU[%d].a=%f\n",i,var_CPU[i].a);
   }

   float pc=2; // IS THIS OK?

   cudaMemcpyToSymbol(c, &pc, sizeof(c)); // IS THIS OK?


   cudaMalloc((void**)&var_GPU,N*sizeof(variable));
   cudaMemcpy(var_GPU,var_CPU,N*sizeof(variable), cudaMemcpyHostToDevice);
   CollisioniGPU<<<10, 10>>>(var_GPU);
   cudaGetLastError();

   cudaMemcpy(var_CPU, var_GPU, N*sizeof(variable),cudaMemcpyDeviceToHost);
   for(i=0;i<N;i++){
      printf("var_CPU[%d].a=%f\n",i,var_CPU[i].a);
   }

   cudaFree(var_GPU);
   free(var_CPU);

   return 0;
}

Here is the kernel.cu:

__device__ void funzione(float *var){

        *var = *var*c;    
}

__global__ void CollisioniGPU(variable *var_GPU){

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

    float a;

    a=var_GPU[id].a;

    funzione(&a);

    var_GPU[id].a = a;      
}

Here is my question: Is this a correct way of using the constant memory? How can I be sure that the multiplication factor c is saved there and not somewhere else?

That is fine. You are basically using the #include to insert code into main.cu — Christian Sarofeen, Mar 11 '15 at 17:07
mmm what do you mean with "#include"? is there something wrong with it? — Federico Gentile, Mar 11 '15 at 17:11
Your code is equivalent to having the code in `kernel.cu` inserted in `main.cu` where you have the line `#include "kernel.cu"` — Christian Sarofeen, Mar 11 '15 at 17:17
possible duplicate of [How to use constant memory for beginners (Cuda C)](http://stackoverflow.com/questions/28987495/how-to-use-constant-memory-for-beginners-cuda-c) — pQB, Mar 12 '15 at 10:34

score 1 · Accepted Answer · answered Mar 12 '15 at 08:22

It seems correct to me, through I would have written things a bit differently.

For definition of c, i would put it in kernel.cu and add extern keyword in main.cu. (And because I am lazy, I often just put these two files in one with two parts divided by some comment block.)

For line with float pc = 2; I often add point and zero, because it's easier to distinguish between integer and floating point variables while reading the code.

And for the cudaMemcpyToSymbol(c, &pc, sizeof(c)); I would have used sizeof(float).

This all is just matter of prefference and clarification of code.

Am I using constant memory correctly?

1 Answers1