shared memory as a parameter of function

Accelerated Computing CUDA CUDA Programming and Performance
1

Dear all;

how can I send a shared memory as a parameter in recursive function. for example this code

__device__ void ex(int *a, int  *t)
{
__shared__ int t[64];

// do some operation 

ex (a,temp);

}

the last line of code gives an error like this:
error : a pointer to shared memory cannot be stored into the parameter buffer
how can i fix that?

2

Define the shared memory variable outside the scope of the ex function. Pass a pointer to it to the ex function.

#include <stdio.h>

__device__ __host__ int fact(volatile int *t){

  int temp = *t;
  if (temp == 2) return temp;
  else {
    *t=temp-1;
    return temp * fact(t);}
}

__global__ void test(int x){

  __shared__ int data[64];
  data[0] = x;
  printf("x = %d, fact(x) = %d\n", x, fact(data));
}

int main(){

  int a = 5;
  printf("cpu fact %d = %d\n", 5, fact(&a));
  test<<<1,1>>>(5);
  cudaDeviceSynchronize();
}
3

thanks

if I want to use (temp) in fact function as follows

temp[tx]=temp[tx]+5;

it gives me error

what can I do?

4

thanks

if I want to use (temp) in fact function as follows

temp[tx]=temp[tx]+5;

it gives me error

what can I do?

5

That’s probably because there is no variable ‘tx’ defined inside the function ‘fact’ that txbob posted above. If by ‘tx’ you mean ‘threadIdx.x’, you can either use that directly, or define your own variable:

int tx = threadIdx.x;
6

I defined int tx = threadIdx.x; as you say, but the problem exixts

7

Show your code, show your compiler invocation, and the resulting error message.

8

You could either do:

int tx = threadIdx.x;

outside your function and pass it as a parameter to the function, or you could define it inside the function. If you define it inside the function, you can’t decorate your function with host anymore, as threadIdx.x is undefined in host code.

Here’s an example of the second approach (defining it inside the recursive function):

#include <stdio.h>

__device__ int fact(volatile int *t){

  int tx = threadIdx.x;
  int temp = t[tx];
  if (temp == 2) return temp;
  else {
    t[tx]=temp-1;
    return temp * fact(t);}
}

__global__ void test(){

  __shared__ int data[64];
  data[threadIdx.x] = threadIdx.x+2;
  printf("x = %d, fact(x) = %d\n", threadIdx.x+2, fact(data));
}

int main(){

  test<<<1,5>>>();
  cudaDeviceSynchronize();
}