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

-

let's have __device__ function:

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

and within __global__ kernel use in way:

uarray = dev_kernel (uarray, params);

where uarray array located in shared memory.

but when use cuda-gdb see addresss of uarray within __global__ kernel get:

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

and within __device__ kernel get:

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

despite error, program in running ok (maybe limitation of cuda-gdb).

so, want know: within __device__ kernel, uarray shared yet? i'm changing array global shared memory , time same (with shared memory time little worse).

so, want know: within __device__ kernel, uarray shared yet?

yes, when pass pointer shared memory device function way, still points same place in shared memory.

in response questions posted below perplexing me, elected show 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;   (int = 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 /home/user2/misc/t249...done. (cuda-gdb) break mykernel breakpoint 1 @ 0x4025dc: file t249.cu, line 9. (cuda-gdb) run starting program: /home/user2/misc/t249 [thread debugging using libthread_db enabled]  breakpoint 1, mykernel () @ t249.cu:9 9       __global__ void mykernel(){ (cuda-gdb) break 14 breakpoint 2 @ 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 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)>>> () @ t249.cu:12 12        (int = 0; i< ssize; i++) (cuda-gdb) continue continuing.  breakpoint 2, mykernel<<<(1,1,1),(1,1,1)>>> () @ t249.cu:14 14        unsigned char *loc = dev_kernel(myshared, 5); (cuda-gdb) print &(myshared[0]) $1 = (@shared unsigned char *) 0x8 ""        ^        |      cuda-gdb telling pointer defined in __shared__ statement, , therefore it's storage implicit , unmodifiable.  (cuda-gdb) print &(u_array) $2 = (@generic unsigned char * @shared *) 0x0        ^                          ^        |                          u_array stored in shared memory.       u_array generic pointer, meaning can point anything.    (cuda-gdb) step dev_kernel(unsigned char * @generic, int) (array_sh=0x1000008 "", params=5)     @ t249.cu:6 6           return array_sh + params; (cuda-gdb) print array_sh $3 = (@generic unsigned char * @register) 0x1000008 ""           ^                      ^           |                    array_sh stored in register.          array_sh generic pointer, can point anything.  (cuda-gdb) print u_array no symbol "u_array" in current context.  (note can't access u_array inside __device__ function, don't understand comment there.)  (cuda-gdb) step mykernel<<<(1,1,1),(1,1,1)>>> () @ 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 stored in shared memory      u_array generic pointer, can point (cuda-gdb)

although haven't provided it, assuming definition of u_array similar mine, based on cuda-gdb output getting.

note indicators @shared not telling kind of memory pointer pointing to, telling either kind of pointer (defined implicitly in __shared__ statement) or else stored (in shared memory).

if doesn't sort out questions, please provide complete example, along complete cuda-gdb session output, have.


Comments

Post a Comment

Popular posts from this blog

c# - How Configure Devart dotConnect for SQLite Code First? -

-
can configure following sample code ? when run following code, error : no entity framework provider found ado.net provider invariant name 'devart.data.sqlite' setting in machine.config <system.data> <dbproviderfactories> <add name="dotconnect sqlite" invariant="devart.data.sqlite" description="devart dotconnect sqlite" type="devart.data.sqlite.sqliteproviderfactory, devart.data.sqlite, version=4.6.287.0, culture=neutral, publickeytoken=09af7300eec23701" /> </dbproviderfactories> </system.data> after add below block code in app.config <?xml version="1.0"?> <configuration>   <configsections> <!-- more information on entity framework configuration, visit http://go.microsoft.com/fwlink/?linkid=237468 --> <section name="entityframework" type="system.data.entity.internal.configfile.entityframeworksection, entit...
Read more

erlang - Saving a digraph to mnesia is hindered because of its side-effects -

-
updated the digraph library in erlang implements data type side-effectful. each operation on data type saves new state in ets table -- instead of returning altered version of itself, more common in functional languages. the problem usage perspective, impedes effort store or pass around state in convenient way requiring me first "collect" state before can start juggle around. the closest solution have seen far serializer/deserializer , these have drawback tied current structure of digraphs instead of operating on abstract type -- prevents future proof solution. update pascal pointed out serializer utilizes interface of digraph, , hence eliminates drawback mentioned above. better, albeit still inconvenient, see no better alternative. what recommendation on how store digraphs? should go different data type altogether? thanks. i not sure understand concern. fact digraph stored in ets transparent in user interface, might want chnge either exchange digr...
Read more

java - Copying object fields -

-
what best way of copying classes when need have 2 independent variables? have simple class: public class myclass { boolean = false; string b= "not empty"; } do need make method : assign(myclass data ) { a= data.a; b= data.b; } is there automatic method copy (duplicate) objects in java? do need make method : pretty close. instead of making method, should make constructor. such constructors called copy constructor , , create them this: myclass(myclass data) { = data.a; b = data.b; } and create copy of instance, use constructor this: myclass obj1 = new myclass(); myclass obj2 = new myclass(obj1); copy constructor can tedious : using copy constructor create deep-copy can tedious, when class has mutable fields. in case, assignment create copy of reference, , not object itself. have create copy of fields (if want deep copy). can go recursive. a better way create deep copy serialize , deserialize object. why not use clone()? ...
Read more