Shared memory address passed to device function is still shared memory?

Question

Let's say I have this __device__ function:

__device__ unsigned char* dev_kernel(unsigned char* array_sh, int params){
    return array_sh + params;
}

And within the __global__ kernel I use it in this way:

uarray = dev_kernel (uarray, params);

Where uarray is an array located in shared memory.

But when i use cuda-gdb to see the addresss of uarray within __global__ kernel I get:

(@generic unsigned char * @shared) 0x1000010 "z\377*"

And within __device__ kernel I get:

(unsigned char * @generic) 0x1000010 <Error reading address 0x1000010: Operation not permitted>

Despite the error, the program in running ok (maybe it is some limitation of cuda-gdb).

So, I want to know: Within the __device__ kernel, uarray is shared yet? I'm changing the array from global to shared memory and the time is almost the same (with shared memory the time is a little worse).

Robert Crovella · Answer 1 · 2013-09-25T03:20:39.533

So, i want to know: Within the __device__ kernel, uarray is shared yet?

Yes, when you pass a pointer to shared memory to a device function this way, it still points to the same place in shared memory.

In response to the questions posted below which are perplexing me, I elected to show a simple example:

$ cat t249.cu
#include <stdio.h>

#define SSIZE 256

__device__ unsigned char* dev_kernel(unsigned char* array_sh, int params){
    return array_sh + params;
}

__global__ void mykernel(){
  __shared__ unsigned char myshared[SSIZE];
  __shared__ unsigned char *u_array;
  for (int i = 0; i< SSIZE; i++)
    myshared[i] = (unsigned char) i;
  unsigned char *loc = dev_kernel(myshared, 5);
  u_array = loc;
  printf("val = %d\n", *loc);
  printf("val = %d\n", *u_array);
}

int main(){

  mykernel<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}
$ nvcc -arch=sm_20 -g -G -o t249 t249.cu
$ cuda-gdb ./t249
NVIDIA (R) CUDA Debugger
5.5 release
....
Reading symbols from /home/user2/misc/t249...done.
(cuda-gdb) break mykernel
Breakpoint 1 at 0x4025dc: file t249.cu, line 9.
(cuda-gdb) run
Starting program: /home/user2/misc/t249
[Thread debugging using libthread_db enabled]

Breakpoint 1, mykernel () at t249.cu:9
9       __global__ void mykernel(){
(cuda-gdb) break 14
Breakpoint 2 at 0x4025e1: file t249.cu, line 14.
(cuda-gdb) continue
Continuing.
[New Thread 0x7ffff725a700 (LWP 26184)]
[Context Create of context 0x67e360 on Device 0]
[Launch of CUDA Kernel 0 (mykernel<<<(1,1,1),(1,1,1)>>>) on Device 0]
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 2, warp 0, lane 0]

Breakpoint 1, mykernel<<<(1,1,1),(1,1,1)>>> () at t249.cu:12
12        for (int i = 0; i< SSIZE; i++)
(cuda-gdb) continue
Continuing.

Breakpoint 2, mykernel<<<(1,1,1),(1,1,1)>>> () at t249.cu:14
14        unsigned char *loc = dev_kernel(myshared, 5);
(cuda-gdb) print &(myshared[0])
$1 = (@shared unsigned char *) 0x8 ""
       ^
       |
     cuda-gdb is telling you that this pointer is defined in a __shared__ statement, and therefore it's storage is implicit and it is unmodifiable.

(cuda-gdb) print &(u_array)
$2 = (@generic unsigned char * @shared *) 0x0
       ^                          ^
       |                          u_array is stored in shared memory.
      u_array is a generic pointer, meaning it can point to anything.  

(cuda-gdb) step
dev_kernel(unsigned char * @generic, int) (array_sh=0x1000008 "", params=5)
    at t249.cu:6
6           return array_sh + params;
(cuda-gdb) print array_sh
$3 = (@generic unsigned char * @register) 0x1000008 ""
          ^                      ^
          |                    array_sh is stored in a register.
         array_sh is a generic pointer, it can point to anything.

(cuda-gdb) print u_array
No symbol "u_array" in current context.
 (note that I can't access u_array from inside the __device__ function, so I don't understand your comment there.)

(cuda-gdb) step
mykernel<<<(1,1,1),(1,1,1)>>> () at t249.cu:15
15        u_array = loc;
(cuda-gdb) step
16        printf("val = %d\n", *loc);
(cuda-gdb) print u_array
$4 = (
    @generic unsigned char * @shared) 0x100000d ......
       ^                         ^
       |                       u_array is stored in shared memory
     u_array is a generic pointer, it can point to anything
(cuda-gdb)

Although you haven't provided it, I am assuming your definition of u_array is similar to mine, based on the cuda-gdb output you are getting.

Note that the indicators like @shared are not telling you what kind of memory a pointer is pointing to, they are telling you either what kind of pointer it is (defined implicitly in a __shared__ statement) or else where it is stored (in shared memory).

If this doesn't sort out your questions, please provide a complete example, along with complete cuda-gdb session output, just as I have.

Thanks, but why cuda-gdb don't say it is @shared just like it did within `__global__` kernel? — Blufter, Sep 24 '13 at 16:45
Are you using standalone cuda-gdb or from within nsight eclipse edition? Which version of CUDA are you using? — Robert Crovella, Sep 24 '13 at 16:55
I'm using standalone cuda-gdb and the version of CUDA is 5.5. — Blufter, Sep 24 '13 at 16:59
Please show the complete definition of `uarray`. I think you are misinterpreting the meaning of the pointer descriptions that are given by cuda-gdb. In fact, it would be best if you showed a complete, compilable, simple example, and show the cuda-gdb command stream and it's output. If you are trying to print `uarray` from within the `__device__` function, it's not in scope there. — Robert Crovella, Sep 25 '13 at 02:48
I've updated my answer with an example cuda-gdb session, decoded for you. — Robert Crovella, Sep 25 '13 at 03:21

Shared memory address passed to device function is still shared memory?

1 Answers1