\n\tmodule simpleOps_m<\/strong><\/span>
\ncontains
\nattributes(global) subroutine inc(a, b)
\nimplicit none
\nreal :: a(:,:)
\nreal, value :: b
\ninteger :: i, j, n(2)
\ni = (blockIdx%x-1)*blockDim%x + threadIdx%x
\nj = (blockIdx%y-1)*blockDim%y + threadIdx%y
\n!threadIdx%x\uff0cthreadIdx%y\u6bcf\u4e2a\u7ebf\u7a0b\u5757\u5185\u7684\u7ebf\u7a0b\u7f16\u53f7\uff08\u4e8c\u7ef4\uff09<\/span>
\n!blockDim%x\uff0cblockDim%y\u6bcf\u4e2a\u7ebf\u7a0b\u5757\u5185\u542b\u6709\u7684\u7ebf\u7a0b\u6570\uff08\u4e8c\u7ef4\uff09<\/span>
\n!blockIdx%x\uff0cblockIdx%y\u7ebf\u7a0b\u5757\u7f16\u53f7\uff08\u4e8c\u7ef4\uff09<\/span>
\nn = size(a)
\nif (i<=n(1) .and. j<=n(2)) a(i,j) = a(i,j) + b
\n!\u4e0d\u80fd\u8d8a\u754c<\/span>
\nend subroutine inc
\nend module simpleOps_m
\n <\/p>\n
program incTest<\/strong><\/span>
\nuse cudafor
\nuse simpleOps_m
\nimplicit none
\ninteger, parameter :: nx=1024, ny=512
\n!\u5b9a\u4e491024\u884c512\u5217\u7684\u77e9\u9635
\nreal :: a(nx,ny), b
\nreal, device :: a_d(nx,ny)
\ntype(dim3) :: grid, tBlock
\na = 1; b = 3
\ntBlock = dim3(32,8,1)
\n!\u5355\u4f4d\u7ebf\u7a0b\u5757\u884c\u8ba1\u7b97\u7ebf\u7a0b\u657032\uff0c\u5355\u4f4d\u7ebf\u7a0b\u5757\u5217\u8ba1\u7b97\u7ebf\u7a0b\u6570\u4e3a8<\/span>
\ngrid = dim3(ceiling(real(nx)\/tBlock%x), ceiling(real(ny)\/tBlock%y), 1)
\n!\u53c2\u4e0e\u884c\u8ba1\u7b97\u7684\u7ebf\u7a0b\u5757\u6570\u5757\u4e3a32\uff0c\u53c2\u4e0e\u5217\u8ba1\u7b97\u7684\u7ebf\u7a0b\u5757\u6570\u4e3a64<\/span>
\n!\u7ebf\u7a0b\u5757\u6570\u4e0d\u80fd\u5927\u4e8emax grid dims<\/span>
\na_d = a
\ncall inc<<<grid,tBlock>>>(a_d, b)
\n!\u8c03\u7528\u683c\u5f0fkernelfunction<<<dimGrid,dimBlock>>>(parameters)
\na = a_d
\nwrite(*,*) 'Max error: ', maxval(abs(a-4))
\nend program incTest<\/p>\n
From\uff1ahttp:\/\/blog.163.com\/zpfzcjndx@126\/blog\/static\/635456812014121104122345\/ <\/span><\/p>\n<\/div>\n","protected":false},"excerpt":{"rendered":"
module simpleOps_m contains attributes(global) subrouti … \u7ee7\u7eed\u9605\u8bfb\u201cCUDA FORTRAN \u6d4b\u8bd5\u7a0b\u5e8f\u201d<\/span><\/a><\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"closed","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[12],"tags":[],"views":1100,"_links":{"self":[{"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/posts\/1475"}],"collection":[{"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/posts"}],"about":[{"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"http:\/\/81.70.49.155\/index.php?rest_route=%2Fwp%2Fv2%2Fcomments&post=1475"}],"version-history":[{"count":3,"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/posts\/1475\/revisions"}],"predecessor-version":[{"id":1478,"href":"http:\/\/81.70.49.155\/index.php?rest_route=\/wp\/v2\/posts\/1475\/revisions\/1478"}],"wp:attachment":[{"href":"http:\/\/81.70.49.155\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=1475"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"http:\/\/81.70.49.155\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=1475"},{"taxonomy":"post_tag","embeddable":true,"href":"http:\/\/81.70.49.155\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=1475"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}