利用CAL进行并行计算

引言：
      对于初学GPGPU的朋友，基本上都是从Brook+，CUDA C或者OpenCL开始的，因为这些接口的Kernel都是基于类C语言进行开发的，相对来说容易很多。就现在而言，CUDA C的接口会成熟一些。但是很多人的显卡是A卡，因为A卡的Graphics相对好一些，对于这些朋友如果想在自己的机器上进行高性能计算，可以选择CAL进行开发。CAL(Compute Abstraction Layer)是AMD的GPGPU的硬件抽象层，它提供了一些底层的功能，从而可以使用户进行高性能的计算。相对前几种接口而言，开发难度可能会大一些，因为CAL的Kernel是用类汇编语言进行编写的，可读性并没有类C语言好。

      本文讲解了一个利用CAL开发的简单的控制台程序。程序的内容非常简单，就是把两个数组中每个元素分别乘起来而已。本文适合对于GPGPU有一定了解的朋友。

正文：
      首先我们看下host端，我们需要做些什么才可以正确设置起来CAL的环境。

            1. 初始化CAL: calInit，调用这个函数就好。

            2. 打开设备。

            3. 创建Context。

      简单的三部设置就可以设置好CAL的环境了。设置好环境后，我们需要编译Kernel代码。也分为三部分：

            1. Compile Kernel的代码（之前需要得到设备的属性，因为CAL会为每个不同的设备编译优化的代码）。

            2. 链接Object。

            3. 读取Module。

      其实Host端会比Brook+复杂一些，但是并没有什么本质的区别，只不过Brook+是封装了这些步骤。所以数据设置部分我就不再重复了，代码里面有的。我们重点来看一下kernel部分吧。

            il_ps_2_0
            dcl_input_interp(linear) v0.xy__
            dcl_resource_id(0)_type(2d,unnorm)_fmtx(float)_fmtx(float)_fmtx(float)_fmtx(float)
            dcl_resource_id(1)_type(2d,unnorm)_fmtx(float)_fmtx(float)_fmtx(float)_fmtx(float)
            dcl_output_generic o0
            flr r0 , v0
            sample_resource(0)_sampler(0) r1 , r0.xyxx
            sample_resource(1)_sampler(1) r2 , r0.xyxx
            mul o0 , r1 , r2
            end

      以上这个简短的程序就是CAL的Kernel部分了，对于刚接触CAL的朋友来说，并不是一眼就能看出来这个程序是做什么的，所以我贴上等同的OpenCL的代码好了。

            __kernel void GPUKernel( const __global float4* data1 ,
                                                 const __global float4* data2 ,
                                                 __global float4* result )
            {
                 uint tidx = get_global_id(0);
                 uint tidy = get_global_id(1);
                 uint dim = get_global_size(0);

                 uint tid= tidy * dim + tidx;

                 result[tid] = data1[tid]+data2[tid];
            }

     上面这个Kernel简单多了，基本没什么东西。我就不多做介绍了。我们重点来看下CAL的那个Kernel。

     第一行的il_ps_2_0是说，语言是IL(Intermediate Language)，ps代表pixel shader，而2_0代表是pixel shader 2.0版本。做通用计算可以用ps还可以用cs(compute shader)，由于前者支持的广泛一些，所以我们这里用ps。

     dcl_input_interp(linear) v0.xy__ ，dcl应该是declera的意思，这里的意思是声明一个变量，这个变量是根据你的thread来确定的，就相当于Opencl里面的get_global_id。这个变量有两点需要注意：

     1. CAL里面的寄存器是没有数据类型的概念的。程序员有责任记住每个变量是什么类型的。这里的索引并不是int型的，而是float型的。这个索引是从(0.5,0.5)开始的，相邻的thread的索引差1.0。

     2. 这个索引是2维的，因为我们外面创建资源的时候创建的是2维的。因为资源的每一维是有长度限制的，在Radeon 4800Series上是8192，5870上是16384。如果我们要处理更多的数据，一维的资源显然不能满足我们的需求，所以我们这里创建了2维的资源。那么我们最大可以创建8192*8192个element，而每个element可以有4个float，那么一共是268435456个数据。这已经是1G的显存了。不过大多数程序不需要一次处理这么大的计算量，所以二维的资源完全可以满足我们的需求。

     下面的两条指令是声明资源，unnorm代表资源里面的数据不是单位化的，后面我们规定了资源里面每个元素的类型。这里我们注意到一个细节，每个元素拥有四个float，事实上amd的GPU在进行计算时，这四路是同时计算的。我们把每四个float分别放置到一个element里面，这样我们可以充分利用AMD GPU的硬件特性进行计算。

     后面就是输出结果的声明。程序的声明部分到这里就已经结束了。下面就是计算的部分了，相对来说更容易看一些。

     flr r0 , v0 指令就是取整指令，这里取整也是四路的，但是只有前两路对我们有意义。

     然后进行采样，就是资源的读取，这里由于资源是二维的，所以我们要输入二维的索引。索引的类型必须为float。

     下面就是最核心也是最简单的计算的部分了，把寄存器r1和r2的内容乘起来而已。这样这个简单的Kernel就结束了。

     其实前面的声明对于刚接触的朋友来说，可能稍微复杂一点，但是熟悉了以后，发现IL也并不是特别难写的。这个程序同样是附有源代码的：
     http://filer.blogbus.com/4730079/resource_47300791259325111j.rar