Declare a global variable to be used as a local one

Question

Is there any way to declare a global variable as if it were to be used as a local one?

I mean:

__device__ int m_myvar;

__global__ void myKernel()
{
  .. do something with m_myvar;
}

I'd like to declare m_myvar as local to the functions where it is used. Any way to do that?

Answer 1

In principle, you can do something like

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

__device__ int m_myvar;

__global__ void myKernel()
{
    m_myvar = threadIdx.x+1;
    printf("myKernel - thread %i - m_myvar %i\n",threadIdx.x,m_myvar);
}

__global__ void myKernel2()
{
    m_myvar = 2*(threadIdx.x+3);
    printf("myKernel2 - thread %i - m_myvar %i\n",threadIdx.x,m_myvar);
}

int main() {

    myKernel<<<1,4>>>();
    myKernel2<<<1,4>>>();
    getch();
    return 0;

}

As @talonmies has observed, operating on m_myvar by different threads in parallel in the two kernels will give rise to race conditions. Therefore, you have to think of using atomic operations when a scalar m_myvar is involved or defining m_myvar as an array as

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

__device__ int m_myvar[4];

__global__ void myKernel()
{
    m_myvar[threadIdx.x] = threadIdx.x+1;
    printf("myKernel - thread %i - m_myvar %i\n",threadIdx.x,m_myvar[threadIdx.x]);
}

__global__ void myKernel2()
{
    m_myvar[threadIdx.x] = 2*(threadIdx.x+3);
    printf("myKernel2 - thread %i - m_myvar %i\n",threadIdx.x,m_myvar[threadIdx.x]);
}

int main() {

    myKernel<<<1,4>>>();
    myKernel2<<<1,4>>>();
    getch();
    return 0;

}