在之前的文章 PGI 介绍了一项旨在弥补NVIDIA的cuda C语言在多核X86架构上的开发计划Using PGI CUDA-x86PGI 已经将CUDA的扩展整合到PGC++® 编译器. 你可以在X86处理器上在正常的PGC++编译器选项中编译CUDA程序使用.cu 后缀:pgcpp -c -O file.cu另外, 您可以启用CUDA的扩展在任何C或C + +源文件- Mcudax86标志：pgcpp -c -O -Mcudax86 file.c
如果选择单独链接，你会需要包括- Mcudax86标志，包括所有适当的CUDA库。它通常会编译器生成x86的对象或可执行文件， __global__或 __device__ 函数属性被编译在一个特殊的模式，使CUDA的内置变量（ threadIdx ， blockIdx等） 。许多内核或设备功能的限制将被删除在x86上，IO的限制等。CUDA运行时API例程，大部分已实施工作在x86主机上。内存管理例程（ cudaMalloc ， cudaFree ， cudaMemcpy ，等等）将分配并复制从主机到主机的数据。因为作为主程序相同的主机上执行的内核，你其实可以简单地传递主机指针的内核，不需要分配到设备存储器和数据复制。我们期待CUDA程序员开发编码，使数据移动电话的条件编译的战略，使移动数据时才会发生针对NVIDIA CUDA设备。这种编码策略，将有助于最大限度地提高多核x86作为CUDA计算，减少或消除虚假数据移动设备的性能。些例程x86主机上没有相应的行为。例如， cudaThreadSetCacheConfig用于设置NVIDIA的费米级GPU的缓存配置; x86主机有没有这样的控制，所以这个调用将CUDA -X86有没有效果。 CUDA - x86的所有操作是同步的，包括内核发射。流管理API例程支持，但管理多个数据流不会产生任何性能上的好处。当pgcpp 编译器编译CUDA-x86扩展时，将定义处理器变量 __PGI_CUDA_X86 ，所以你可以写你的程序与特定的GPU和x86的具体行为。PGDBG ®调试已得到增强，以及支持的CUDA- x86程序。看到有关调试与PGDBG CUDA -X86在此PGInsider问题的姊妹篇。CUDA -X86支持所有x86平台上， 32位和64位，目前PGI编译器支持包括Linux， Mac OS X和Windows的。CUDA-x86 语言限制PGC + +编译器是CUDA -X86已验证的代码样本集的NVIDIA CUDA SDK提供改良，修改目前的执行情况，我们在这里讨论有关的限制。Driver API: 当前实现只支持CUDA运行时API ，这包括对流行的（ <<< ... >>> ）的语法，和CUDA的支持...... API例程。正在探索对较低级别的驱动程序API的支持。Textures and Surfaces: 这些大多是用在GPU上采取的2维和3维纹理内存高速缓存的优势。纹理和表面的支持仍在实施，将在未来的版本。但他们将在x86处理器上，没有什么特别的性能优势。CUDA Arrays: 这些是专门分配的数据，用于支持纹理和表面构造，将在同一时间支持。OpenGL and DirectX Interoperability: OpenGL和DirectX的互操作性的好处是明确的GPU ， CUDA代码和图形共享相同的内存和硬件。目前CUDA - x86实现不包括这些API例程，并讨论未来的支持下。Thrust: Thrust库尚未移植到CUDA -X86Warp-synchronous Programming: 当两个线程在同一个线程块需要通过共享内存通信，一般需要包括__ syncthreads （ ）之间的通话由一个线程共享内存的分配和使用由其他线程，内存，一个简单的例子，将是一个循环数组的值求和。在这个内核，每个线程计算当地的总和，在第一循环迭代通过数组的大小为N ，这是线程块中的线程数块。每个线程，然后将其存储到共享存储器阵列与当地的总和，第二个循环做了从N部分和二进制减少一个总和。#define N 128__global__ void sumit( float* a, int m ){ int i, n; float localsum = 0.0f; __shared__ float S[N]; for( i = threadIdx.x; i < m; i += N ) localsum += a; i = threadIdx.x; S = localsum; __syncthreads(); /* all threads done with localsum */ n = N; while( (n>>1) > 0 ){ /* add S[0:n-1] += S[n:2*n-1] */ if( i < n ) S += S[i+n]; __syncthreads(); /* all partial sums done */ } /* here, S[0] has the final sum */}如果需要同步的线程数量是经线的大小相等，我们不需要__ syncthreads ， （）在第二个循环调用。对于GPU ，经大小是32 ，所以一些程序员将以前的内核优化如下：#define N 128__global__ void sumit( float* a, int m ){ int i, n; float localsum = 0.0f; __shared__ float S[N]; for( i = threadIdx.x; i < m; i += N ) localsum += a; i = threadIdx.x; S = localsum; __syncthreads(); /* all threads done with localsum */ n = N; while( (n>>1) > 32 ){ /* add S[0:n-1] += S[n:2*n-1] */ if( i < n ) S += S[i+n]; __syncthreads(); /* all partial sums done */ } while( (n>>1) > 0 ){ /* add S[0:n-1] += S[n:2*n-1] */ if( i < n ) S += S[i+n]; /* the last 32 elements are all done by one warp * so we don't need to synchronize */ } /* here, S[0] has the final sum */}CUDA -X86没有弥补在扭曲的线程，因此经同步规划将无法正常工作。值内置在warpSize变量， CUDA -X86 cudaGetDeviceProperties经大小返回的值是1 。如果您使用wrap同步，你应该对测试warpSize ，而不是字面常量32。CUDA-x86 Performance此版本的编译器使用就是我们所说的仿真模式。编译器的PGI CUDA Fortran的仿真模式，创建一个任务，每个CUDA线程和一台主机系统上的每个核心的工作线程。 工作者线程的时间在一个单一的线程块执行的所有任务（ CUDA线程）的份额。工作者线程将执行一个单一的任务，直到达到一个同步点（ __syncthreads ） ，然后切换到一个不同的任务，直到所有任务完成。线程块执行一次;多核并行使用一个块内执行多个线程。这一个面向功能的释放，而不是一个绩效导向的释放。虽然有充足的并行线程块内保持一个多核心的忙，你写了一个程序为GPU的方式一般是完全错误的，以这种方式执行多核心方案。为了得到接合内存负载和GPU上的商店，你组织你的数据，使相邻的线程，在线程块（ threadIdx.x连续值）访问设备内存在相邻的位置。与此CUDA -X86的版本，这意味着不同的内核将访问内存中的相邻位置，这意味着他们将有众多的缓存冲突。在仿真模式下， CUDA程序将运行明显比NVIDIA的GPU相同的程序慢。今年晚些时候，我们将推出一个优化的实施，将改变并行模型，并用矢量单一核心上执行多个线程并行。我们期望的性能将显着改善，但它仍然可能会是一个因素，比GPU慢，这是不奇怪，因为在高端GPU的大规模并行处理。另一个性能指标是比较CUDA程序到本机OpenMP实现相同的算法。 CUDA程序总是会有额外的开销，内存管理和矩形螺纹块和网格处理，但我们的目标是向本地OpenMP并行使用CUDA代码的性能。
SummaryPGI已发布的第一个本地的CUDA的C / c + +多核心x86处理器的执行。这第一个版本实现跨线程块内的线程，使用多核心并行CUDA运行时API 。缺少一些功能，如纹理，将在未来几个月内，和一个高度优化的代码生成器也正在开发。
