CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第六节

原文链接

第六节:全局内存和CUDA RPOFILER 

Rob Farber 是西北太平洋国家实验室(Pacific Northwest National Laboratory)的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究,并且是几个新创企业的合伙人。大家可以发邮件到[email protected]与他沟通和交流。

在关于CUDA(Compute Unified DeviceArchitecture,即计算统一设备架构的简称)的系列文章的第二节,我探讨了内存性能和在reverseArray_multiblock_fast.cu.内使用共享内存。在本节,我探讨使用CUDA PROFILER检测全局内存。

本系列文章的细心读者已经了解了第四节和第五节里讨论的两个反向数组样例,但是仍然困扰他们的是,共享内存版本为什么比全局内存版的速度要快一些。回忆下共享内存版本吧,reverseArray_multiblock_fast.cu,内核将数组数据从全局内存复制到共享内存,然后回到全局内存,而较慢的内核reverseArray_multiblock.cu,仅将数据从全局内存复制到全局内存。因为全局内存性能比共享内存要慢100~150倍,慢得多的全局存储器性能占据了两个示例的绝大部分运行时。为什么共享存储器版本更快?回答这个问题需要先了解更多有关全局内存的信息,还需要使用来自CUDA开发环境的附加工具--特别是CUDAPROFILER。CUDA软件的配置简单快捷,因为文本和可视化版本的profiler都在CUDA启动的设备上读取硬件配置计数器。启动文本配置非常简易:设置开始和控制profiler的环境变量。使用可视化profiler同样很简单:启动cudaprof并开始在GUI中进行单击操作。通过配置可以了解许多有价值的信息。配置事件集合完全由支持CUDA的设备内部的硬件来处理。然而,经过配置的内核不再具有异步特征。只有在每个内核完成之后,才将结果报告给主机,这样可以最小化所有通信带来的影响。

全局内存
了解如何有效使用全局内存是成为一名CUDA编程高手的基本要求。下面对全局内存进行了简要介绍,应该可以有助于你了解reverseArray_multiblock.cu和reverseArray_multiblock_fast.cu之间的性能区别。如有需要,以后的专栏文章会继续探索如何有效利用全局内存。同时,我们会采用图示的方式详细探讨全局内存(见CUDA Programming Guide的第5.1.2.1节)。

只有当全局存储器访问能够合并到一个half-warp时,硬件才能以最少的事务量获取(或存储)数据,全局存储器才能交付最高的存储器带宽。CUDAComputeCapability设备(1.0和1.1)能够在单个64字节或128字节事务中获取数据。如果无法合并存储器事务,那么将会为half-warp中的每个线程发出一个独立的存储器事务,这不是期望的结果。未合并的存储器操作的性能损失取决于数据类型的大小。CUDA文档对各种数据类型大小决定的预期性能降低给出了一些简单指南:

  • 32位数据类型将减慢大约10x
  • 64位数据类型将减慢大约4x
  • 128位数据类型将减慢大约2x

当满足下列条件时,数据块的half-warp中的所有线程执行的全局存储器访问可以被合并到G80架构上一个有效的存储器事务中:
线程访问32、64或128位数据类型。

事务的所有16个字所在的分段的大小必须和内存事务大小一致(当为128位字时,为内存事务大小的2倍)。这就意味着起始地址和校准非常的重要了。

线程必须依次访问这些字:half-warp中的第k个线程必须访问第k个字。注意:不是warp中的所有线程都需要访问某个线程所访问的存储器才能进行合并。这称为发散warp。

较新的架构(比如GT200系列设备)的合并要求比刚才讨论的架构更宽松。我们将在未来的专栏中更深入地讨论它们之间的架构差异。从本专栏的主题看,可以肯定,如果经过调优的代码能够在支持CUDA的G80设备上进行很好的合并,那么它将能够在GT200设备上进行很好地合并。

启动和控制文本配置

控制CUDA profiler的文本版本的环境变量是:

  • CUDA_PROFILE – 设置为1(或0)可以启用(或禁用)profiler
  • CUDA_PROFILE_LOG – 设置为日志文件的名称(默认设置为./cuda_profile.log)
  • CUDA_PROFILE_CSV – 设置为1(或0)可以启用(或禁用)使用逗号分隔的日志版本。
  • CUDA_PROFILE_CONFIG – 指定最多带有4个信号的配置文件

最后一点非常重要,因为一次只能配置四个信号。通过在名为CUDA_PROFILE_CONFIG的文件中的单独行上指定名称,开发人员可以用profiler收集以下任何事件:gld_incoherent:未合并的全局存储器负载单元的数量

  • gld_coherent:已合并的全局存储器负载单元的数量
  • gst_incoherent:未合并的全局存储器存储单元的数量
  • gst_coherent:已合并的全局存储器存储单元的数量
  • local_load:局部存储器负载单元的数量
  • local_store:局部存储器存储单元的数量
  • branch:线程执行的分支事件的数量
  • divergent_branch:warp中发散分支的数量
  • instructions:指令计数
  • warp_serialize:warp中基于与共享或常量存储器的地址冲突进行序列化的线程数量
  • cta_launched:执行的线程块

Profiler 计数器注意问题:
注意,性能计数器值与单个的线程活动无联系。实际上,这些值代表了线程warp内的事件。例如,一个线程warp中的一个不连贯的存储将会递增gst_incoherent一次。因此,最终的计数器值存储的是关于所有warp中的所有不连贯存储的信息。

此外,profiler仅以GPU中的一个多处理器为目标,因此计数器值与为特定内核启动的warp的总数不会相关。因此,当使用profiler内的性能计数器选项时,用户应该总是启动足够的线程块以确保为目标多处理器分配固定百分比的工作。实际上,NVIDIA建议最好启动至少100个块,以获得一致的结果。

结果就是,用户不应该期待计数器值与通过检查内核代码所确定的数值一致。计数器值最好用于确定未优化和已优化代码之间的相对性能差异。例如,如果profiler报告软件的初始部分有一定数量的未合并全局负载,那么很容易确定更精细的代码版本是否会利用更少数量的未合并负载。在大多数情形下,我们的目标是将未合并的全局负载数量减少为0,因此,计数器值对于跟踪此目标的实现进度非常有用。

配置结果
我们来使用profiler.看下reverseArray_multiblock.cu 和reverseArray_multiblock_fast.cu。在本样例中,我们会在Linux下的bash shell中对环境变量和配置文件进行如下设置:

1 export CUDA_PROFILE=1
2 export CUDA_PROFILE_CONFIG=$HOME/.cuda_profile_configexport CUDA_PROFILE=1
3 export CUDA_PROFILE_CONFIG=$HOME/.cuda_profile_config

在Linux下使用bash比较Profile配置和环境变量

 1 gld_coherent
 2 gld_incoherent
 3 gst_coherent
 4 gst_incoherent
 5 [code]
 6
 7 CUDA_PROFILE_CONFIG文件内容
 8 运行reverseArray_multiblock.cu可执行文件,在./cuda_profile.log中生成以下配置报告:
 9 [code]
10 method,gputime,cputime,occupancy,gld_incoherent,gld_coherent,gst_incoherent,gst_coherent
11 method=[ memcopy ]
12 gputime=[ 438.432 ]
13 method=[ _Z17reverseArrayBlockPiS_ ]
14 gputime=[ 267.520 ]
15 cputime=[ 297.000 ]
16 occupancy=[ 1.000 ]
17 gld_incoherent=[ 0 ]
18 gld_coherent=[ 1952 ]
19 gst_incoherent=[ 62464 ]
20 gst_coherent=[ 0 ]
21 method=[ memcopy ]
22 gputime=[ 349.344 ] 

reverseArray_multiblock.cu配置报告
类似地,运行reverseArray_multiblock_fast.cu可执行文件生成以下输出,这些输出会覆盖.cuda_profile.log中以前的输出

 1 method,gputime,cputime,occupancy,gld_incoherent,gld_coherent,gst_incoherent,gst_coherent
 2 method=[ memcopy ]
 3 gputime=[ 449.600 ]
 4 method=[ _Z17reverseArrayBlockPiS_ ]
 5 gputime=[ 50.464 ]
 6 cputime=[ 108.000 ]
 7 occupancy=[ 1.000 ]
 8 gld_incoherent=[ 0 ]
 9 gld_coherent=[ 2032 ]
10 gst_incoherent=[ 0 ]
11 gst_coherent=[ 8128 ]
12 method=[ memcopy ]
13 gputime=[ 509.984 ]

reverseArray_multiblock_fast.cu配置报告

比较这两个profiler结果,可看到reverseArray_multiblock_fast.cu内没有不连贯的存储,而 reverseArray_multiblock.cu却相反,它包含很多不连贯存储。看一下reverseArray_multiblock.cu的源,并看一下您是否可以修复不连贯存储的性能问题。修复之后,测量一下这两个程序彼此的相对速度。

为方便起见,列表1显示了reverseArray_multiblock.cu的情况,列表2显示了reverseArray_multiblock_fast.cu的情况。

 1 // includes, system
 2 #include <stdio.h>
 3 #include <assert.h>
 4 // Simple utility function to check for CUDA runtime errors
 5 void checkCUDAError(const char* msg);
 6 // Part3: implement the kernel
 7 __global__ void reverseArrayBlock(int *d_out, int *d_in)
 8 {
 9     int inOffset  = blockDim.x * blockIdx.x;
10     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
11     int in  = inOffset + threadIdx.x;
12     int out = outOffset + (blockDim.x - 1 - threadIdx.x);
13     d_out[out] = d_in[in];
14 }
15 ////////////////////////////////////////////////////////////////////////////////
16 // Program main
17 ////////////////////////////////////////////////////////////////////////////////
18 int main( int argc, char** argv)
19 {
20     // pointer for host memory and size
21     int *h_a;
22     int dimA = 256 * 1024; // 256K elements (1MB total)
23     // pointer for device memory
24     int *d_b, *d_a;
25     // define grid and block size
26     int numThreadsPerBlock = 256;
27     // Part 1: compute number of blocks needed based on array size and desired block size
28     int numBlocks = dimA / numThreadsPerBlock;
29     // allocate host and device memory
30     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
31     h_a = (int *) malloc(memSize);
32     cudaMalloc( (void **) &d_a, memSize );
33     cudaMalloc( (void **) &d_b, memSize );
34     // Initialize input array on host
35     for (int i = 0; i < dimA; ++i)
36     {
37         h_a[i] = i;
38     }
39     // Copy host array to device array
40     cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
41     // launch kernel
42     dim3 dimGrid(numBlocks);
43     dim3 dimBlock(numThreadsPerBlock);
44     reverseArrayBlock<<< dimGrid, dimBlock >>>( d_b, d_a );
45     // block until the device has completed
46     cudaThreadSynchronize();
47     // check if kernel execution generated an error
48     // Check for any CUDA errors
49     checkCUDAError("kernel invocation");
50     // device to host copy
51     cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
52     // Check for any CUDA errors
53     checkCUDAError("memcpy");
54     // verify the data returned to the host is correct
55     for (int i = 0; i < dimA; i++)
56     {
57         assert(h_a[i] == dimA - 1 - i );
58     }
59     // free device memory
60     cudaFree(d_a);
61     cudaFree(d_b);
62     // free host memory
63     free(h_a);
64     // If the program makes it this far, then the results are correct and
65     // there are no run-time errors.  Good work!
66     printf("Correct!\n");
67
68     return 0;
69 }
70 void checkCUDAError(const char *msg)
71 {
72     cudaError_t err = cudaGetLastError();
73     if( cudaSuccess != err)
74     {
75         fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
76         exit(EXIT_FAILURE);
77     }
78 }

reverseArray_multiblock.cu

 1 // includes, system
 2 #include <stdio.h>
 3 #include <assert.h>
 4 // Simple utility function to check for CUDA runtime errors
 5 void checkCUDAError(const char* msg);
 6 // Part 2 of 2: implement the fast kernel using shared memory
 7 __global__ void reverseArrayBlock(int *d_out, int *d_in)
 8 {
 9     extern __shared__ int s_data[];
10     int inOffset  = blockDim.x * blockIdx.x;
11     int in  = inOffset + threadIdx.x;
12     // Load one element per thread from device memory and store it
13     // *in reversed order* into temporary shared memory
14     s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];
15     // Block until all threads in the block have written their data to shared mem
16     __syncthreads();
17     // write the data from shared memory in forward order,
18     // but to the reversed block offset as before
19     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
20     int out = outOffset + threadIdx.x;
21     d_out[out] = s_data[threadIdx.x];
22 }
23 ////////////////////////////////////////////////////////////////////////////////
24 // Program main
25 ////////////////////////////////////////////////////////////////////////////////
26 int main( int argc, char** argv)
27 {
28     // pointer for host memory and size
29     int *h_a;
30     int dimA = 256 * 1024; // 256K elements (1MB total)
31     // pointer for device memory
32     int *d_b, *d_a;
33     // define grid and block size
34     int numThreadsPerBlock = 256;
35     // Compute number of blocks needed based on array size and desired block size
36     int numBlocks = dimA / numThreadsPerBlock;
37     // Part 1 of 2: Compute the number of bytes of shared memory needed
38     // This is used in the kernel invocation below
39     int sharedMemSize = numThreadsPerBlock * sizeof(int);
40     // allocate host and device memory
41     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
42     h_a = (int *) malloc(memSize);
43     cudaMalloc( (void **) &d_a, memSize );
44     cudaMalloc( (void **) &d_b, memSize );
45     // Initialize input array on host
46     for (int i = 0; i < dimA; ++i)
47     {
48         h_a[i] = i;
49     }
50     // Copy host array to device array
51     cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
52     // launch kernel
53     dim3 dimGrid(numBlocks);
54     dim3 dimBlock(numThreadsPerBlock);
55     reverseArrayBlock<<< dimGrid, dimBlock, sharedMemSize >>>( d_b, d_a );
56     // block until the device has completed
57     cudaThreadSynchronize();
58     // check if kernel execution generated an error
59     // Check for any CUDA errors
60     checkCUDAError("kernel invocation");
61     // device to host copy
62     cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
63     // Check for any CUDA errors
64     checkCUDAError("memcpy");
65     // verify the data returned to the host is correct
66     for (int i = 0; i < dimA; i++)
67     {
68         assert(h_a[i] == dimA - 1 - i );
69     }
70     // free device memory
71     cudaFree(d_a);
72     cudaFree(d_b);
73     // free host memory
74     free(h_a);
75     // If the program makes it this far, then the results are correct and
76     // there are no run-time errors.  Good work!
77     printf("Correct!\n");
78     return 0;
79 }
80
81 void checkCUDAError(const char *msg)
82 {
83     cudaError_t err = cudaGetLastError();
84     if( cudaSuccess != err)
85     {
86         fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
87         exit(EXIT_FAILURE);
88     }
89 }

reverseArray_multiblock_fast .cu

时间: 2024-10-10 10:24:21

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第六节的相关文章

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第七节

第七节:使用下一代CUDA硬件,快乐加速度 原文链接 Rob Farber 是西北太平洋国家实验室(Pacific Northwest National Laboratory)的高级科研人员.他在多个国家级的实验室进行大型并行运算的研究,并且是几个新创企业的合伙人.大家可以发邮件到[email protected]与他沟通和交流. 在关于CUDA(Compute Unified Device Architecture,即计算统一设备架构的简称)的系列文章的第六节,我探讨了使用CUDA PROFI

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第十节

原文链接 第十节:CUDPP, 强大的数据平行CUDA库Rob Farber 是西北太平洋国家实验室(Pacific Northwest National Laboratory)的高级科研人员.他在多个国家级的实验室进行大型并行运算的研究,并且是几个新创企业的合伙人.大家可以发邮件到[email protected]与他沟通和交流. 在关于CUDA(Compute Unified Device Architecture,即计算统一设备架构的简称)的系列文章的第9节,我探讨了如何使用CUDA拓展高

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第一节

原文链接 第一节 CUDA 让你可以一边使用熟悉的编程概念,一边开发可在GPU上运行的软件. Rob Farber 是西北太平洋国家实验室(Pacific Northwest National Laboratory)的高级科研人员.他在多个国家级的实验室进行大型并行运算的研究,并且是几个新创企业的合伙人.大家可以发邮件到[email protected]与他沟通和交流. 您是否有兴趣在使用高级语言(比如C编程语言)编程时,通过标准多核处理器将性能提升几个数量级?您是否期待拥有跨多个设备的伸缩能力

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第九节

原文链接 第九节:使用CUDA拓展高等级语言 Rob Farber 是西北太平洋国家实验室(Pacific Northwest National Laboratory)的高级科研人员.他在多个国家级的实验室进行大型并行运算的研究,并且是几个新创企业的合伙人.大家可以发邮件到[email protected]与他沟通和交流. 在关于CUDA(Compute Unified Device Architecture,即计算统一设备架构的简称)的系列文章的第八节,我探讨了使用CUDA利用库.在本小节,我

CUDA:Supercomputing for the Masses (用于大量数据的超级计算)-第三节

原文链接 恭喜!通过对CUDA(Compute Unified DeviceArchitecture,即计算统一设备架构的首字母缩写)系列文章第一节和第二节,您现在已经是能够使用CUDA的程序员了,您可以创建和运行在支持CUDA的设备上使用成百上千同步线程的程序.在第二节的incrementArrays.cu中,我提供了一个常见的CUDA应用程序模式的工作样例——将数据移动到设备,运行一个或多个内核以进行计算并获得结果.本质上,只需使用您自己的内核并加载自己的数据(我在本篇专栏文章的示例中就是这

为数据挖掘小组写的一个用于造数据的小程序

最近有个数据挖掘的项目,要求在文本里面写入随机字母并且要1000W个 于是就写了个程序用来造数据并记录一下 程序写的时候遇到的问题 1 未考虑内存溢出的情况,大批量的把数据写入导致内存溢出 以后需要谨慎对待 目前完整版 package test; import java.io.File; import java.io.FileWriter; import java.io.IOException; import org.apache.commons.io.FileUtils; import org

大数据架构和模式(四)——了解用于大数据解决方案的原子模式和复合模式

摘要:本文中介绍的模式有助于定义大数据解决方案的参数.本文将介绍最常见的和经常发生的大数据问题以及它们的解决方案.原子模式描述了使用.处理.访问和存储大数据的典型方法.复合模式由原子模式组成,并根据大数据解决方案的范围进行分类.由于每个复合模式都有若干个维度,所以每个模式都有许多变化.复合模式使得业务和技术用户可以应用一个结构化方法为大数据问题建立范围,并定义高级的解决方案. 简介 本系列的 第 3 部分 介绍了大数据解决方案的逻辑层.这些层定义了各种组件,并对它们进行分类,这些组件必须处理某个

Memblaze 联手PMC推出用于超大规模数据中心的高性能PCIe SSD

Memblaze 联手PMC 推出用于超大规模数据中心的高性能PCIe SSD Memblaze 产品在容量.灵活度及延迟等方面均领先业界       引领大数据连接.传送以及存储,提供创新半导体及软件解决方案的PMC?公司(纳斯达克代码:PMCS)今日宣布,忆恒创源科技有限公司( Memblaze Technology Co. Ltd)在其下一代PCIe产品中采用了PMC的Flashtec? NVM Express (NVMe) NVMe控制器. 基于Flashtec的Memblaze的PBl

Druid:一个用于大数据实时处理的开源分布式系统

Druid是一个用于大数据实时查询和分析的高容错.高性能开源分布式系统,旨在快速处理大规模的数据,并能够实现快速查询和分析.尤其是当发生代码部署.机器故障以及其他产品系统遇到宕机等情况时,Druid仍能够保持100%正常运行.创建Druid的最初意图主要是为了解决查询延迟问题,当时试图使用Hadoop来实现交互式查询分析,但是很难满足实时分析的需要.而Druid提供了以交互方式访问数据的能力,并权衡了查询的灵活性和性能而采取了特殊的存储格式. Druid功能介于PowerDrill和Dremel