GPU编程基础：CUDA Fortran实践指南

5星 · 超过95%的资源需积分: 49 83 浏览量更新于2024-07-19 1 收藏 4.67MB PDF 举报

"CUDA Fortran for Scientists and Engineers" 是一本专为使用Fortran进行科学计算和工程应用开发的专家和工程师编写的指南，旨在教授如何利用CUDA Fortran来加速GPU并行计算。这本书由Gregory Ruetsch和Massimiliano Fatica共同撰写，来自NVIDIA Corporation。它提供CUDA Fortran的基础知识以及一些典型的示例，帮助读者在不将CUDA Fortran编程本身变成目标的情况下，有效地掌握GPU编程。 CUDA（Compute Unified Device Architecture）是NVIDIA推出的一种并行计算平台和编程模型，主要用于提升图形处理器（GPU）的计算性能。CUDA Fortran是CUDA的一个扩展，允许开发者使用Fortran语言来编写高性能计算程序，充分利用GPU的并行处理能力，尤其适合于科学计算、数值模拟等领域。本书的核心内容可能包括以下几个方面： 1. **CUDA基础**：介绍CUDA的架构，包括GPU的工作原理、多线程模型（如线程块和网格）、内存层次（全局内存、共享内存、常量内存和纹理内存）以及数据传输机制。 2. **CUDA Fortran语法**：讲解如何在Fortran代码中嵌入CUDA语句，如定义设备变量、使用CUDA函数、进行内核调用等。 3. **并行编程概念**：讨论并行计算的基本概念，如并行度、同步、原子操作和流，以及如何优化并行性能。 4. **性能优化**：提供最佳实践，指导如何编写高效、低延迟的CUDA Fortran代码，包括如何减少数据传输、最大化内存带宽利用率、优化线程调度等。 5. **实例分析**：书中会包含多个实际的科学计算或工程问题案例，演示如何使用CUDA Fortran解决这些问题，并进行性能比较。 6. **错误处理与调试**：介绍CUDA编程中常见的错误类型，以及如何调试CUDA Fortran代码，确保其正确性和效率。 7. **工具和资源**：可能会提及CUDA开发工具，如NVIDIA的Nsight工具链，以及社区和文档资源，帮助读者进一步学习和提高。 8. **许可和版权**：书籍明确指出未经许可，不得复制或以任何形式传播书中的内容，同时也提供了获取权限的信息。通过这本书，读者可以系统地学习CUDA Fortran编程，掌握GPU并行计算的关键技术和策略，从而加速现有的Fortran应用，提高计算效率。对于那些在科研或工程领域需要高效计算的人来说，这是一份宝贵的资源。

12 CHAPTER 1 Introduction

In both of the CUDA Fortran examples we have discussed, the kernel code accesses the x ﬁelds of

the predeﬁned variables, and as you might expect, these data types can accommodate multidimensional

arrays, which we explore next.

1.3.3 Multidimensional arrays

We can easily extend our example to work on a multidimensional array. This is facilitated since the

predeﬁned variables in device code are of a derived type dim3, which contains x, y, and z ﬁelds. In

terms of the host code, thus far we have speciﬁed the blocks per grid and threads per block execution

conﬁguration parameters as integers, but these parameters can also be of type dim3. Using other ﬁelds

of the dim3 type, the multidimensional version of our code becomes:

1 module simpleOps_m

2 contains

3 attributes(global) subroutine increment(a, b)

4 implicit none

5 integer :: a(:,:)

6 integer , value :: b

7 integer :: i, j, n(2)

9 i = (blockIdx%x-1)* blockDim%x + threadIdx%x

10 j = (blockIdx%y-1)* blockDim%y + threadIdx%y

11 n(1) = size(a ,1)

12 n(2) = size(a ,2)

13 if (i <=n (1) .and. j <=n (2)) a(i,j) = a(i,j) + b

14 end subroutine increment

15 end module simpleOps_m

19 program incrementTest

20 use cudafor

21 use simpleOps_m

22 implicit none

23 integer , parameter :: nx=1024 , ny =512

24 integer :: a(nx,ny), b

25 integer , device :: a_d(nx ,ny)

26 type(dim3 ) :: grid , tBlock

28 a=1

29 b=3

31 tBlock = dim3 (32,8,1)

32 grid = dim3 (ceiling(real(nx)/ tBlock%x), &

33 ceiling(real(ny)/ tBlock%y), 1)

34 a_d = a

1.4 Determining CUDA hardware features and limits 13

35 call increment <<<grid , tBlock >>>(a_d , b)

36 a = a_d

38 if (any(a /= 4)) then

39 write (* ,*) ’**** Program Failed **** ’

40 else

41 write (*,*) ’Program Passed ’

42 endif

43 end program incrementTest

After declaring the parameters nx and ny along with the host and device arrays for this two-

dimensional example, we declare two variables of type dim3 used in the execution conﬁguration on

line 26. On line 31 the three components of the dim3 type specifying the number of threads per

block are set; in this case each block has a 32 × 8 arrangement of threads. In the following two lines,

the ceiling function is used to determine the number of blocks in the x and y dimensions required

to increment all the elements of the array. The kernel is then launched with these variables as the

execution conﬁguration parameters in line 35. In the kernel code, the dummy argument a is declared

as a two-dimensional array and the variable n as a two-element array, which, on lines 11 and 12, is

set to hold the size of a in each dimension. An additional index j is assigned a value on line 10 in an

analogous manner to i on line 9, and both i and j are checked for in-bound access before a(i,j) is

incremented.

1.4 Determining CUDA hardware features and limits

There are many different CUDA-capable devices available, spanning different product lines (GeForce

and Quadro

as well as Tesla) in addition to different generations of architecture. We have already

discussed the limitation of the number of threads per block, which is 1024 on Kepler and Fermi-based

hardware and 512 for earlier architectures, and there are many other features and limits that vary among

devices. In this section we cover the device management API, which contains routines for determining

the number and types of CUDA-capable cards available on a particular system and what features and

limits such cards have.

Before we go into the device management API, we should brieﬂy discuss the notion of compute

capability. The compute capability of a CUDA-enabled device indicates the architecture and is given

in Major.Minor format. The Major component of the compute capability reﬂects the generation of the

architecture, and the Minor component reﬂects the revision within that generation. The very ﬁrst CUDA-

enabled cards were of compute capability 1.0. Fermi-generation cards have compute capabilities of 2.x,

and Kepler-generation cards have compute capabilities of 3.x. Some f eatures of CUDA correlate with

the compute capability; for example, double precision is available with cards of compute capability 1.3

and higher. Other features do not correlate with compute capability but can be determined through the

device management API.

The device management API has routines for getting information on the number of cards avail-

able on a system as well as for selecting a card from among available cards. This API makes use of

1.4 Determining CUDA hardware features and limits 15

44 write (*,"(’ Max Grid Dims: ’,2(i0,’ x ’),i0)") &

45 prop% maxGridSize

46 write(*,"(’ Max Block Dims: ’,2(i0 ,’ x ’),i0)") &

47 prop% maxThreadsDim

48 write(*,"(’ Max Threads per Block: ’,i0 ,/)") &

49 prop% maxThreadsPerBlock

51 enddo

53 end program deviceQuery

This code determines the number of CUDA-capable devices attached to the system from the

cudaGetDeviceCount() routine on line 10 and then loops over each device, retrieving the

device properties from the cudaGetDeviceProperties() routine. This code lists only a small

portion of the ﬁelds available in the cudaDeviceProp type. A full list of the members of the

cudaDeviceProp derived type is provided in the CUDA Toolkit Reference Manual, available online

from NVIDIA.

We list the output of this code on a variety of Tesla devices of different compute capabilities. The

ﬁrst Tesla device, released in 2007, was the Tesla C870, with a compute capability of 1.0, and on a

system with this device we obtain the following result:

 

One CUDA device found

Device Number: 0

Device Name: Tesla C870

Compute Capability: 1.0

Number of Multiprocessors: 16

Max Threads per Multiprocessor: 768

Global Memory (GB): 1.500

Execution Configuration Limits

Max Grid Dims: 65535 x 65535 x 1

Max Block Dims: 512 x 512 x 64

Max Threads per Block: 512



Note that the enumeration of devices is zero-based rather than unit-based.TheMax Threads per

Multiprocessor number refers to the maximum number of concurrent threads that can reside

on a multiprocessor. The Global Memory indicated in the following line of output is the amount

of available memory in device DRAM. The ﬁrst two lines under Execution Configuration

Limits denote the limits in each dimension of the ﬁrst two execution conﬁguration parameters:

the number and conﬁguration of thread blocks in a kernel launch, and the number and conﬁguration

of threads in a thread block. Note that for this compute capability, grids must be a two-dimensional

剩余315页未读，继续阅读

feng-yr

粉丝: 4
资源: 3

GPU编程基础：CUDA Fortran实践指南

Beginning Julia Programming For Engineers and Scientists 无水印原版pdf

Fortran for Scientists & Engineers [4th Edition, 2017]

Fortran for Scientists and Engineers(4th) 无水印pdf

Fortran90-95 for Scientists and Engineers

Fortran 95,2003 for scientists and engineers

fortran 90 for engineers and scientists

Fortran 95/2003 Scientists and Engineers 书中代码

fortran_for_engineers

Essential MATLAB for Engineers and Scientists - 7th.pdf

SVM_CudaFortran:在Cuda Fortran中实现支持向量机

最新资源