cuda fortran

CUDA Fortran 共享内存

共享内存shared memory的使用

1. shared memory 因为放在芯片上，所以访问速度比local memory快的多
2. 每一个thread block 中的shared memory是动态分配的，thread block中的所有线程读取一样的shared memroy.Thread block中的线程可以读取另一个线程从global memory 装载到shared memory中的数据

module reverse_m
    !实现一个一维数组反转
    implicit none
    integer,device::n_d
    
    contains
    
    attributes(global) subroutine staticReverse(d)
        implicit none
        real::d(:)
        integer::t,tr
        real,shared::s(64) !在编译的时候就声明了shared memory的大小
        
        !原始和反转的索引计算
        t=threadidx%x      
        tr=side(d)-t+1
    
        s(t)=d(t)   !数据从global memory 复制到shared memory
        !每一个线程读取shared memory 中由另一个线程写入的数据，
        !需要确保所有线程完成数据导入的工作
        !global memroy中的读写操作都是通过索引T的，反索引tr用来读取shared memroy
        call syncthreads()
    !Thread block 中的一个线程可以读取另一个线程从global memory 装载到shared memory中的数据
    !因此shared memory 被应用于global memory联合访问行不通的地方
        d(t)=s(tr)
    end subroutine staticReverse

    attributes(global) subroutine dynamicReversel(d)
        real::d(:)
        integer::t,tr
        real,shared::s(*) !假定大小的方式，其大小隐性由启动kernel的第三个参数决定
    
        t=threadidx%x
        tr=size(d)-t+1
        
        s(t)=d(t)
        call syncthreads()
        d(t)=s(tr)
    end subroutine dynamicReversel

    attributes(global) subroutine dynamicReversel2(d,nsize)
        real::d(nsize)
        integer,value::nsize
        integer::t,tr
        real,shared::s(nsize)
        
        t=threadidx%x
        tr=nsize-t+1
        
        s(t)=d(t)
        call sycnthreads()
        d(t)=s(tr)
    end subroutine dynamicReversel2
    
   attributes(global) subroutine dynamicReversel3(d)
        real::d(n_d)
        real,shared::s(n_d)
        integer::t,tr
        
        t=thredidx%x
        tr=n_d-t+1
        
        s(t)=d(t)
        call sycnthreads()
        d(t)=s(tr)
    end subroutine dynamicReversel2
    
    program sharedExample
        use cudafor
        use reverse_m
        implicit none
        
        integer,parameter::n=64
        real::a(n),b(n),c(n)
        real,device::d_d(n)
        
        type(dim3)::grid,tblock
        integer::i
        
        tblock=dim3(n,1,1)
        grid=dim3(1,1,1)
        
        do i=1,n
            a(i)=i
            r(i)=n-i+1
        end do
        
        d_d=a
        call staticReverse<<<grid,tblock>>>(d_d)
        d=d_d
        write(*,*)"static case max error",maxval(abs(r-d))
        
        call dynamicReversel<<<grid,tblock,4*threadBlock%x>>>(d_d)
        d=d_d
        write(*,*)"dynamic case 1 max error",maxval(abs(r-d))
        
        call dynamicReversel2<<<grid,tblock,4*threadBlock%x>>>(d_d,n)
        d=d_d
        write(*,*)"dynamic case 2 max error",maxval(abs(r-d))
        
        n_d=n
        d_d=a
        call dynamicReversel3<<<grid,threadBlock,4*threadBlock%x>>>(d_d)
        d=d_d
        write(*,*)"dynamic case 3 max error",maxval(abs(r-d))
        
    end program sharedExample

pgfortran -Mcuda sharedExample.f90
sharedExample.exe

使用多个动态分配shared memory数组，特别是想使用不同类型的shared memory，那么就使用dynamicreverse2和3，如果使用假定形状，编译器不知道这样的动态数组分配多大的shared memory

shared memory存储冲突

为了同时访问达到高的带宽，shared memory被分成大小相等的快，

一个内存区域的多个访问需要映射到一个memory bank.访问需要依次执行，硬件将有bank冲突的一个内存需求分解成多个独立、无冲突的内存需求，一个half warp中的所有的线程读取shared memory同样的位置，这样会导致数据广播

在shared memory中，连续的32bit字节映射到连续的bank，每一个bank每一个时钟周期都是32bit大小带宽，shared memory 带宽是每个时钟周期32bit,

32bit=4B

integer(kind=4),其中4为4byte

bit 位 计算机表示数据最小的单位

byte 字节 1byte=8bit

一个字符=2个字节B

1k=1024B

计算力为1.x的设备，wrap大小为32，bank的数量为16.

对于计算力2.0的设备，wrap大小为32，bank的数量是32，一个warp需求的shared memory不会像计算力为1.x的设备拆成两个，意味冲突不会出现在两个half warp彼此的线程中