[CUDA FORTRAN] Problem with SHARED attribute

Hello,

I’ve written a simple 2D reduction code in order to fin the minimum value of a 16*16 matrix. I found a problem with the “shared” attribute: in particular, if I use a “device” array the code performs correctly, instead with the “shared” attribute it doesn’t works (the result is displayed as a row of stars).

I’ve uploaded the CUDA FORTRAN project on MediaFire:

http://www.mediafire.com/?h1qoz5g8velrcte

Note: to perform the 2D reduction with the shared memory feature, just uncomment the call the the second kernel and comment the call to the first one (using device memory).


Thanks in advance,

Nicola

Hi Nicola,

You’re using dynamic shared memory (shared automatic arrays) but failed to specify the size of the memory in the kernel launch. Hence, your kernel is crashing and ‘res’ is returning garbage.

To fix, either change your shared arrays to be fixed size (i.e. change dim_mx_x and dim_mx_y to be parameters) or pass in the shared memory size. For example:

call simple_red_shrd  1,dimBlock_s_mx,((dim_mx_x*dim_mx_y*pr)+(dim_mx_x*2*pr))   (dim_mx_x,dim_mx_y,s_mx_dev,res)

Hope this helps,
Mat

Thanks for you help, I chose this way and fixed the problem. In order to help other people who may read this thread, what I needed to do was the following:

  • Declare the 2D array’s dimensions (dim_mx_x, dim_mx_y) as parameters in the variables module:
integer,parameter :: dim_mx_x = 16,dim_mx_y = 16
  • Remove the dimensions from the kernel call:
call simple_red_shrd<<<1>>>(s_mx_dev,res)
  • Inside the kernel, declare the dimensions with a “USE ONLY” statement:
use simple_variables,only: dim_mx_x,dim_mx_y

Nicola