CUDA my shared memory code not working, what am I missing?

Question

I am attempting to implement dynamic shared memory and it is not working. Please review the code and tell me what am I missing - the problem seems to be around gpu_configuration().

The following is a basic dynamic shared memory code. I have compared it with NVIDIA-developer-blog/code-samples/.../shared-memory/shared-memory.cu and can't figure out what is missing.

If I were to delete the function gpu_configuration(), it works fine but I get an illegal memory access with the function gpu_configuration there. I use this function as part of another piece of code and everything works fine there.

I use a Quadro 2000 card on Kubuntu 14.4, CUDA 7.0 - the card details are printed out by gpu_configuration and are listed below.

BTW, the shared-memory.cu code works fine on my machine so it is not a problem with the card or with Shared Memory.

shmem.cu:

#include <stdio.h>
inline void gpuAssert (cudaError_t code, const char *file, const char *func, int line);
typedef struct gpu_config {
    int     n_threads;  // execution of kernel with given index
    int     n_blocks;   // bundle of threads - also warp size
    int     n_grid;     // bundle of blocks
    int     dev_count;  // number of cuda devices
    size_t  shmem;      // sh_mem per block
    size_t  free_mem;   // free memory on the card
    size_t  tot_mem;    // total memory on the card
    struct  cudaDeviceProp  dev_prop;   // device properties
} gpu_config;

#define CUDA_GLOBAL_CHECK {gpuErrChk (cudaPeekAtLastError ()); gpuErrChk (cudaDeviceSynchronize ());}
#define gpuErrChk(ans)      {gpuAssert((ans), \
__FILE__, __func__, __LINE__);}
#define CudaDbgPrn(M, ...)  {printf ("DevDEBUG:%s:%s:%d: " M "\n", \
    __FILE__, __func__, (int) __LINE__, ##__VA_ARGS__);}
#define Dbg(M, ...)     fprintf(stderr, "DEBUG %s:%s:%d: " M "\n", __FILE__, \
    __func__, __LINE__, ##__VA_ARGS__)
inline void gpuAssert (cudaError_t code, const char *file, const char *func, int line)
{
    if (code != cudaSuccess) {
        fprintf(stderr,"CUDA call from file:%s func:%s %d: %s:%s failed\n", file, func, line, cudaGetErrorName(code), cudaGetErrorString(code));
        exit (code);
    }
}

static void gpu_configuration (gpu_config *gc);
static void gpu_configuration (gpu_config *gc)
{
    int i = 0;
    gpuErrChk (cudaDeviceReset ());     // reset device

    gpuErrChk (cudaGetDeviceCount (&gc -> dev_count));
    Dbg("Device count %d", gc -> dev_count);

    gpuErrChk (cudaGetDeviceProperties (&(gc -> dev_prop), i));
    gc -> n_threads = gc -> dev_prop.maxThreadsPerBlock;
    gc -> n_blocks = gc -> dev_prop.warpSize;
    dim3 block (gc -> n_blocks);
    gc -> n_grid = (gc -> n_blocks + block.x - 1) / block.x;

    gc -> shmem = gc -> dev_prop.sharedMemPerBlock;
    gpuErrChk (cudaMemGetInfo (&(gc -> free_mem), &(gc -> tot_mem)));

    Dbg ("Dev prop name: %s, tot_mem: %u sharedMemPerBlock %u\nwarpSize %d maxThreadsPerBlock %d\nmaxthreads per mprocessor %d",
    gc -> dev_prop.name, (unsigned) gc -> dev_prop.totalGlobalMem,
    (unsigned) gc -> dev_prop.sharedMemPerBlock,
    gc -> dev_prop.warpSize, gc -> dev_prop.maxThreadsPerBlock,
    gc -> dev_prop.maxThreadsPerMultiProcessor);
}

#define         MAX_SIZE        4000
#define         NUM             2
// #define          NUM         32

__global__ void kernel(int *d_data)
{
    extern __shared__ int sdata[];

    sdata[threadIdx.x] = threadIdx.x;
    __syncthreads ();

    CudaDbgPrn ("sdata [%u]=%u", (unsigned) threadIdx.x, (unsigned) sdata[threadIdx.x]);
    CudaDbgPrn ("d_data [%u]=%u", (unsigned) threadIdx.x, (unsigned) d_data[threadIdx.x]);
    d_data[threadIdx.x] = sdata[threadIdx.x];

    CudaDbgPrn ("sdata [%u]=%u d_data [%u]=%u", (unsigned) threadIdx.x, (unsigned) sdata[threadIdx.x], (unsigned) threadIdx.x, (unsigned) d_data[threadIdx.x]);
}

int main()
{
    int *d_data;
    gpuErrChk (cudaMalloc ((void**)&d_data, sizeof(int) * MAX_SIZE));
    gpuErrChk (cudaMemset (d_data, '\0', sizeof(int) * MAX_SIZE));

    gpu_config gc;
    gpu_configuration (&gc);

    kernel<<<1, NUM, (NUM * sizeof (int))>>> (d_data);
    CUDA_GLOBAL_CHECK;

    cudaFree(d_data);
    return 0;
}

And here's what I get on the command line:

rinka@Desktop:~/Documents/dev/code$ nvcc -Xptxas -v shmem_test.cu -o shmem
ptxas info    : 139 bytes gmem, 40 bytes cmem[14]
ptxas info    : Compiling entry function '_Z6kernelPi' for 'sm_20'
ptxas info    : Function properties for _Z6kernelPi
    40 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 21 registers, 40 bytes cmem[0]

rinka@Desktop:~/Documents/dev/code$ ./shmem
DEBUG shmem_test.cu:gpu_configuration:36: Device count 1
DEBUG shmem_test.cu:gpu_configuration:51: Dev prop name: Quadro 2000, tot_mem: 1073414144 sharedMemPerBlock 49152
warpSize 32 maxThreadsPerBlock 1024
maxthreads per mprocessor 1536
DevDEBUG:shmem_test.cu:kernel:65: sdata [0]=0
DevDEBUG:shmem_test.cu:kernel:65: sdata [1]=1
CUDA call from file:shmem_test.cu func:main 82: cudaErrorIllegalAddress:an illegal memory access was encountered failed

Answer 1

When you call gpu_configuration(&gc), the cudaDeviceReset() call inside it deallocates all of the previously allocated memory on the current device. Therefore, d_data becomes invalid and causes the kernel to fail.

You may remove the cudaDeviceReset() call to fix the issue. Alternatively, the gpu_configuration call should be the first function call in your program so that subsequent memory allocations remain valid.