# Recover shared memory used for parameter passing

**URL:** <https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730>\
**Category:** CUDA Programming and Performance\
**Created:** [October 30, 2007, 4:05pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730 "2007-10-30T16:05:41Z")\
**Posts on this page:** 12\
**Page:** 1

<div class="post-metadata">

**Author:** ![jpape](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jpape](https://forums.developer.nvidia.com/u/jpape)\
**Post date:** [October 30, 2007, 4:05pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/1 "2007-10-30T16:05:41Z")

</div>

I’m developing an application that would optimally use all 16K of shared memory for data, but since CUDA passes **global** function arguments via the shared memory, the full 16K isn’t available.

Is there any way to recover the shared memory used for argument passing after copying the arguments to thread registers?

---

<div class="post-metadata">

**Author:** ![sphyraena](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sphyraena](https://forums.developer.nvidia.com/u/sphyraena)\
**Post date:** [October 30, 2007, 4:13pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/2 "2007-10-30T16:13:37Z")

</div>

> [@](#):
>
> I’m developing an application that would optimally use all 16K of shared memory for data, but since CUDA passes **global** function arguments via the shared memory, the full 16K isn’t available.
> 
> Is there any way to recover the shared memory used for argument passing after copying the arguments to thread registers?
> 
> [snapback]272347[/snapback]

How about if you simply passed all parameters via constant memory? Then your kernel will have no arguments.

---

<div class="post-metadata">

**Author:** ![jpape](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jpape](https://forums.developer.nvidia.com/u/jpape)\
**Post date:** [October 30, 2007, 4:16pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/3 "2007-10-30T16:16:51Z")

</div>

> [@](#):
>
> How about if you simply passed all parameters via constant memory? Then your kernel will have no arguments.
> 
> [snapback]272350[/snapback]

That sounds like a good option.

Is there a performance penalty in writing the values to constant memory and accessing them vs. CUDA putting the parameter values in shared memory before the kernel starts? If so, what would the penalty be?

---

<div class="post-metadata">

**Author:** ![asadafag](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/asadafag/32/216303_2.png) [@asadafag](https://forums.developer.nvidia.com/u/asadafag)\
**Post date:** [October 30, 2007, 4:22pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/4 "2007-10-30T16:22:41Z")

</div>

You can’t use 16k shared mem even if you don’t have parameter. Some is reserved for blockIdx, threadIdx and stuff  
There shouldn’t be much penalty of passing via const unless there’re divergently-read look-up tables.

---

<div class="post-metadata">

**Author:** ![wumpus](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/wumpus/32/12040_2.png) [@wumpus](https://forums.developer.nvidia.com/u/wumpus)\
**Post date:** [October 30, 2007, 9:20pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/5 "2007-10-30T21:20:57Z")

</div>

asadafag: I’m not entirely sure yet if these count towards the 16k limit or are in their own private memory space

---

<div class="post-metadata">

**Author:** ![sphyraena](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sphyraena](https://forums.developer.nvidia.com/u/sphyraena)\
**Post date:** [October 30, 2007, 9:38pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/6 "2007-10-30T21:38:30Z")

</div>

CUDA adds a 16-byte shared memory overhead to all kernels. But I’m not sure what it contains. My guess is that threadIdx, blockIdx and other such variables are stored in registers, not shared memory.

---

<div class="post-metadata">

**Author:** ![wumpus](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/wumpus/32/12040_2.png) [@wumpus](https://forums.developer.nvidia.com/u/wumpus)\
**Post date:** [October 30, 2007, 10:11pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/7 "2007-10-30T22:11:49Z")

</div>

In disassembled cubin it appears those parameters (except for threadIdx, which is passed in register r0) are in a separately addressed piece of shared memory, which is read-only. It is completely separate from the rw memory we call ‘shared memory’

```auto
0x0: "%gridflags", # lower u16 is gridid

0x1: "%ntid.x",    # checked

0x2: "%ntid.y",

0x3: "%ntid.z",

0x4: "%nctaid.x",

0x5: "%nctaid.y",

0x6: "%ctaid.x",

0x7: "%ctaid.y",

0x8: "%ctaid.z",  # extrapolated

0x9: "%nctaid.y",  # ptx ISA

```

Parameters start at offset 0x10 (4\*0x4) of the writable shared memory area. Other declared shared memory variables are immediatly after that. As to what is at offset 0x00 I don’t know. These are indeed 16 bytes overhead. I should try reading them out some time and see if they contain the same as the %blockIdx registers.

---

<div class="post-metadata">

**Author:** ![asadafag](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/asadafag/32/216303_2.png) [@asadafag](https://forums.developer.nvidia.com/u/asadafag)\
**Post date:** [October 31, 2007, 2:58am UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/8 "2007-10-31T02:58:52Z")

</div>

I once tried to use 16k for scan, and it indeed failed. A simple test confirmed it:

```auto
#include <stdio.h>

__global__ void aaa(int *a){

	a[0]=7777;

}

int main(){

	int *a,b;

	cudaMalloc((void**)&a,4);

	cudaMemset(a,0,4);

	aaa<<<1,1,16384>>>(a);

	cudaMemcpy(&b,a,4,cudaMemcpyDeviceToHost);

	printf("%d\n",b);

	return 0;

}

```

The kernel launch failed.

---

<div class="post-metadata">

**Author:** ![sphyraena](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sphyraena](https://forums.developer.nvidia.com/u/sphyraena)\
**Post date:** [October 31, 2007, 3:23am UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/9 "2007-10-31T03:23:39Z")

</div>

> [@](#):
>
> Parameters start at offset 0x10 (4\*0x4) of the writable shared memory area. Other declared shared memory variables are immediatly after that. As to what is at offset 0x00 I don’t know. These are indeed 16 bytes overhead. I should try reading them out some time and see if they contain the same as the %blockIdx registers.
> 
> [snapback]272510[/snapback]

Fascinating. What variables are in 0xa, 0xb, …, 0xf?

---

<div class="post-metadata">

**Author:** ![wumpus](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/wumpus/32/12040_2.png) [@wumpus](https://forums.developer.nvidia.com/u/wumpus)\
**Post date:** [October 31, 2007, 8:15am UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/10 "2007-10-31T08:15:17Z")

</div>

Seems I was completely wrong about my separate registers, the block parameters are in shared memory just like anything else, and thereby, writable too:

```auto
#include <cuda_runtime_api.h>

#include <cuda.h>

#include <algorithm>

__global__ void my_kernel(uint16_t *data)

{

extern __shared__ uint16_t x[];

    x[-16] = 0x1234;

    __syncthreads();

    data[0] = x[-16]; // %gridflags

    data[1] = x[-15]; // %ntid.x

    data[2] = x[-14]; // %ntid.y

    data[3] = x[-13]; // %ntid.z

    data[4] = x[-12]; // %nctaid.x

    data[5] = x[-11]; // %nctaid.y

    data[6] = x[-10]; // %ctaid.x

    data[7] = x[-9];  // %ctaid.y

}

int main()

{

    int width = 8;

    int size = width*4;

   uint16_t *data, *gdata;

   cudaMalloc((void**)&gdata, size);

    data = (uint16_t*)malloc(size);  

   for(int x=0; x<8; ++x)

    {

        dim3 block_size(1,2,1);

        dim3 grid_size(8,8,1); 

        int shared_size = 0;  

       my_kernel<<<grid_size, block_size, shared_size>>>(gdata);

       cudaMemcpy((void*)data, (void*)gdata, size, cudaMemcpyDeviceToHost);

       for(int x=0; x<width; ++x)

            printf("%04x ", data[x]);

        printf("\n");

    }

   return 0;

}

```

So… yes, the first 16 (8\*2) bytes of shared memory are the block parameters, and you can never get rid of that overhead. You can overwrite them though if you feel really lucky :P (by indexing into negative shared memory)

Parameter 0x9, 0xA, 0xB etc don’t exist, those would be the normal (user parameters).

Edit: I updated decude to reflect accordingly

---

<div class="post-metadata">

**Author:** ![prkipfer](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@prkipfer](https://forums.developer.nvidia.com/u/prkipfer)\
**Post date:** [October 31, 2007, 1:06pm UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/11 "2007-10-31T13:06:37Z")

</div>

> [@](#):
>
> ```auto
>    data[1] = x[-15]; // %ntid.x
> 
> ```

Hihi, nice one [External Media](http://hqnveipbwb20/public/style_emoticons/<#EMO_DIR#>/w00t.gif "Media hosted on another site. Click to open in a new tab.")

Peter

---

<div class="post-metadata">

**Author:** ![James\_Malcolm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@James\_Malcolm](https://forums.developer.nvidia.com/u/James_Malcolm)\
**Post date:** [July 25, 2008, 12:02am UTC](https://forums.developer.nvidia.com/t/recover-shared-memory-used-for-parameter-passing/1730/12 "2008-07-25T00:02:46Z")

</div>

> [@](#):
>
> Fascinating. What variables are in 0xa, 0xb, …, 0xf?

I played around with this some and I don’t think these are missing parameters. I think that the compiler puts the dynamic shared memory after the kernel parameters and then aligns it to the next 16-byte boundary. Since you have one parameter, it skips three words and then starts the dynamic shared segment.