首页 计算机科学

# 什么是异构并行计算

计算系统中同时包括多种不同架构的处理器

  • X86+GPU、X86+FPGA、ARM/Power+ GPU 等

  • 不同架构处理器对应不同的指令集

异构计算环境本身要求应用程序支持并行计算

  • 无论是 X86 处理器，还是 ARM 和 GPU 处理器以及 DSP， 这所有的处理器都是多核向量处理器，要发挥多种处理器混合平台的性能也必须要采用并行的编程方式

# 为什么需要异构计算

  • 不同处理器架构适合处理不同类型的计算问题
    • 将不同类型的计算合理分布到不同类型硬件上能够获得 更好的计算性能，比如将逻辑判别计算调度到 X86 架构的通用 CPU，而将需要能够大规模向量并行的计算部分调度到 GPU。
  • 为特定应用设计更优架构的处理器
    • 处理器架构设计可以依据应用的具体特点来优化，可兼顾功耗与性能。
    • 但性能、功耗与通用性不可兼得，需要根据实际情况合理取舍。

# 列举五种常见的异构计算平台

  • C++ AMP —— Microsoft

  • CUDA —— Nvidia

  • OpenCL —— Khronos Group

  • OpenACC —— Cray , CAPS , Nvidia, PGI

  • OpenGL —— Khronos Group

  • RenderScript —— Google

# 简述异构并行计算的分类方法

异构计算可分为系统异构计算（shc-system heterogeneous computing）和网络异构计算（nhc-network heterogeneous computing）两大类

SHC 以单机多处理器形式提供多种计算类型，而 NHC 则以网络连接的多计算机形式提供多种计算类型

  1. 系统异构计算（SHC）

    分为单机多计算方式和单机混合计算方式两大类

    • 前者在同一时刻允许以多种计算方式执行任务
    • 后者在同一时刻只允许以一种计算方式执行任务，但在不同时刻计算可从一种方式自动切换到另一种方式，如 simd (single instruction multiple data) 和 mimd ( multiple instruction multiple data) 方式间的切换。
    • 高性能计算（GPU 等）、半定制计算芯片（FPGA）、专用计算芯片（ASIC）、神经网络计算（NPU）

  2. 网络异构计算（NHC）

    分为同类异型多机方式和异类混合多机方式两类

    • 同类异型多机方式中所使用的多机，它们的结构属同一类，即支持同一种并行性类型（如 simd、mimd、向量等类型之一），但型号可能不同，因此性能可以各有差异。通常的 now (Net of Workstations, 工作站网络) 或 cow (Cluster of Workstations, 工作站集群) 为同类同型多机方式，因此可看成是同类异型多机方式中的特例。
    • 异类混合多机方式中所使用的多机，它们的结构则属不同类型。
    • 分布式计算、集群计算、网格计算、云计算

# 简述 GPU 适合的应用场景以及不合适的场景

GPU 适用场景

GPU 只有在计算高度数据并行任务时才能发挥作用。在这类任务中，需要处理大量的数据，数据的储存形式类似于规则的网格，而对这写数据的进行的处理则基本相同。这类数据并行问题的经典例子有：图像处理，物理模型模拟 （如计算流体力学），工程和金融模拟与分析，搜索，排序

更适用于大规模规则化向量化的计算

GPU 不适用场景

  • 需要复杂数据结构的计算如树，相关矩阵，链表，空间细分结构等，则不适用于使用 GPU 进行计算。

  • 串行和事务性处理较多的程序

  • 并行规模很小的应用，如只有数个并行线程

  • 需要 ms 量级实时性的程序

  • 不适用于逻辑复杂的计算

# 简述一下 CPU、GPU、FPGA 和 ASIC 的优缺点

CPU：通用性、适用复杂计算；灵活、易用、通用；性能较低

GPU：适用批量数据并行计算；高性能、高功耗

FPGA：使用不规则数据进行并行计算；性能好、能效比高、灵活

AISC：适用数据并行计算；高性能、低功耗；专用电路不可修改

# 简述 PTX 编程语言

PTX（Parallel Thread eXecution）作用类似于汇编，是为动态编译器设计的输入指令序列。这样，不同的显卡使用不同的机器语言，而动态编译器却可以运行相同的 PTX。这样做使 PTX 成为了一个稳定的接口 —— 后向兼容性，更长的寿命，更好的可拓展性和更高的性能。但在一定程度上也限制了工程的自由发挥。这种技术保证了兼容性，但也使新一代的产品必须拥有上代产品的所有能力，这样才能让今天的 PTX 代码在未来的系统上仍然可以运行。

CUBIN、C

# 简述 CUDA 异构计算过程

可用于计算的硬件分为两部分

  • Host (主机，CPU)，负责控制 / 指挥 GPU 工作
  • Device（GPU），协处理

Host+Device 异构并行 C 应用程序

  • Host 端串行 C 代码
  • Device 端 SPMD 并行化 kernel C 代码

# 简述 CUDA 函数都有哪些限制

  1. 硬件限制：线程分配，和 block 设置要满足协处理器的硬件参数
  2. 内存容量限制：GPU 内存容量有限，对于大规模数据处理任务，可能会出现内存不足的情况。你需要合理管理内存使用，避免过度分配内存。
  3. 并行同步限制：对于一些计算，数据需要同步，需要处理好主机和协处理器的数据传输

线程分组：线程分组对应的维度必须能被整除，分组的大小不能超过 1024。

假设 GPU 线程是一维的，共 8 个，则可以选择每 2 个 GPU 线程为 1 组或者每 4 个 GPU 线程为 1 组，但不能选择每 3 个 GPU 线程为 1 组，因为剩下 2 个 GPU 线程不足 1 组。

Warp 是硬件特性带来的概念，在 CUDA C 语言中是透明的（除 vote 函数），但应用中不能忽略

一个 warp 中有 32 个线程，这是因为 SM 中有 8 个 SP，执行一条指令的延迟是 4 个周期，使用了流水线技术

一个 half warp 中有 16 个线程，这是因为执行单元的频率是其他单元的两倍，每两个周期才进行一次数据传输

SIMT 编程模型: SIMT 是对 SIMD（Single Instruction, Multiple Data，单指令多数据）的一种变形。

两者的区别在于：SIMD 的向量宽度是显式的，固定的，数据必须打包成向量才能进行处理

而 SIMT 中，执行宽度则全由硬件自动处理了.(每个 block 中的 thread 数量不一 定是 32)

SIMT 中的 warp 中的每个线程的寄存器都是私有的，它们只能通过 shared memory 来进行通信

# 简述 GPU 的 Share memory 的 Bank Conflict 问题？

寄存器与local memory

对每个线程来说，寄存器都是线程私有的 —— 这与 CPU 中一样。如果寄存器被消耗完，数据将被存储在本地存储器 (local memory)。 Local memory 对每个线程也是私有的。线程的输入和中间输出变量将被保存在寄存器或者本地存储器中.

Shared Memory Bank 为了获得高带宽，shared memory 被分为了 16（或 32）个等大小内存 块 (banks)，单位是 32-bit。

相邻数据在不同 bank 中，对 16 （或 32）余数相同的数据在同一 bank

Shared memory 是用于线程间通信的共享存储器。共享存储器是一块可以被同一 block 中的所有 thread 访问的可读写存储器

访问共享存储器几乎和访问寄存器一样快，是实现线程间通信的延迟最小的方法。

bank : 堆

如果 warp 访问 shared memory，对于每个 bank 只访问不多于一个内存地址，那么只需要一次内存传输就可以了，否则需要多次传输，因此会降低内存带宽的使用。

bank 间是并行内的，bank 内是串的 (并非不能并，而是出于对关键资源的保护)

Warp 有三种典型的获取 shared memory 的模式：

  • Parallel access：最通常的模式，多个地址分散在多个 bank。这个模式一般暗示，一些（也可能是全部）地址请求能够被一次传输解决。理想情况是，获取无 conflict 的 shared memory 的时，每个地址都落在不同的 bank 中。
  • Serial access：最坏的模式，多个地址落在同一个 bank。如果 warp 中的 32 个 thread 都访问了同一个 bank 中的不同位置，那就是 32 次单独的请求，而不是同时访问了。
  • Broadcast access：一个地址读操作落在一个 bank。也是只执行一次传输，然后传输结果会广播给所有发出请求的 thread。这样的话就会导致带宽利用率低。

以 warp 访问 shared memory 为例，当多个地址请求同时发生时，如果这些地址落在同一个 bank 中，就会产生冲突。这种冲突会导致需要多次传输，降低内存带宽的使用效率。而理想的无 conflict 情况是每个地址都落在不同的 bank 中，这样可以通过一次传输解决多个地址请求，提高内存访问效率。

stride 为奇数时才不会存在 bank conflicts

# 简述什么是可分页内存和分页锁定内存，如何进行资源分配？

可分页内存：这是通过操作系统 API（如 C 语言中的 malloc() 、C++ 中的 new() ）在主机上分配的存储空间

分页锁定内存：是由 CUDA 函数（如 cudaHostAlloc() ）在主机内存上分配的。其重要属性是主机的操作系统不会对这块内存进行分页和交换操作，确保该内存始终驻留在物理内存中，不会被存储到虚拟内存

# 什么是统一寻址，其优势是啥？

在 host 内存与 device 显存之间根据访问需要自动迁移数据，同时保证 host 和 device 都可访问，应用程序并不需要知道访问时数据所在的具体位置

需要进行显式同步（以保证前一步骤中数据更新操作全部完成）

优势：编程简化，数据访问性能提升

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMallocManaged(&ret, 1000 * sizeof(int));
    AplusB<<<1, 1000>>>(ret, 10, 100);
    cudaDeviceSynchronize(); // Synchronize, wait for all kernels to finish
    for (int i = 0; i < 1000; i++) {
        printf("%d: A+B = %d\n", i, ret[i]);
    }
    cudaFree(ret);
    return 0;
}
__device__ __managed__ int ret[1000];
__global__ void AplusB(int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    AplusB<<<1, 1000>>>(10, 100);
    cudaDeviceSynchronize(); // Synchronize
    for (int i = 0; i < 1000; i++) {
        printf("%d: A+B = %d\n", i, ret[i]);
    }
    return 0;
}
# CUDA 程序体系结构分哪几类，他们之间的关系是啥？
  • CUDA 函数库（CUDA Library）
  • CUDA 运行时 API（Runtime API）
  • CUDA 驱动 API（Driver API）
三级结构
CUDA 应用程序可以通过直接调用底层的 CUDA 驱动来调用 GPU 硬件进行并行计算，也可以使用对底层驱动进行了一定封装的 CUDA 运行时库来简化编程过程
对一些标准的数学函数计算，也可以直接调用 CUFFT 、 CUBLAS 以及 CUDPP 等标准函数库进一步简化编程过程。
# 如何识别一个函数到底属于驱动级别的还是运行时级别的？
驱动 API 是一种基于句柄的底层接口，大多数对象通过句柄被引用，其函数前缀均为 cu
在调用驱动 API 前必须进行初始化，然后创建一个 CUDA 上下文，该上下文关联到特定设备并成为主机线程的当前上下文
运行时 API 对驱动 API 进行了一定的封装，隐藏了其部分实现细节，因此使用起来更为方便，简化了编程的过程，因此我们实际中更多使用的是运行时 API. 运行时 API 不需要初始化，在调用运行时 API 时，会自动初始化。
因此只需要观察是否是 cu 开头，和是否需要手动初始化就能判断一个函数是否是驱动级别的，反之是运行时级别的（ 对 CUDA 函数
bool initCUDA() {
    if (culnit(0) != CUDA_SUCCESS) return false; // 初始化 CUDA 运行时
    int deviceCount = 0;
    cuDeviceGetCount(&deviceCount); // 获取当前系统中的设备数量
    if (deviceCount == 0) {
        fprintf(stderr, "There is no device supporting CUDA.\n");
        return false;
    }
    CUdevice cuDevice = 0;
    cuDevicesGet(&cuDevice, 0); // 获取 0 号设备对象的句柄
    CUcontext cuContext;
    cuCtxCreate(&cuContext, 0, cuDevice); // 创建 CUDA 上下文
    //... 其他代码
    return true;
}
运行时 API 初始化
bool initCUDA() {
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0) {
        fprintf(stderr, "There is no device supporting CUDA.\n");
        return false;
    }
    cudaSetDevice(0); // ****
    //... 其他代码
    return true;
}
# 请列举 5 种常见的 CUDA 函数库，简述他们各自适合的场景？
  • cuBLAS 函数库：
    cuBLAS（CUDA Basic Linear Algebra Subprograms）库是一个利用 CUDA 进 行加速版本的完整标准矩阵与向量的运算库
    主要特性完全支持全部 152 个标准 BLAS 例程，支持单精度、 双精度、复数和二重复数等数据类型，支持多 GPU 与并发内核
    Fortan 语言
  • cuSPARSE 函数库
    包含了一系列处理稀疏矩阵的基本的线性代数子程式
    在稠密向量格式和稀疏矩阵向量格式之间的操作、在稀疏矩阵格式和稠密向量格式之间的操作、在稀疏矩阵格式和一组稠密向量之间的操作
  • cuDPP 函数库
    以高效使用 GPU 进行基本的数据并行运算，如并行前缀求和（扫描），并行排序（数字）以及稀疏矩阵－向量相乘算法等
  • cuFFT 函数库
    CUFFT（CUDA Fast Fourier Transform）库是一个利用 CUDA 进行傅里叶变换的函数库
    主要特性是可以实现复数与实数的 1 维、2 维和 3 维变换，其中 1 维变换最大为 1.28 亿个元素，同时数据布局灵活，允许在单个元素与阵列维度之间使用任意幅度
  • NPP 函数库
    高性能 GPU 处理，专注于图像、视频、信号处理领域
  • cuRAND – Random Number Generation (RNG) Library
# 为什么需要进行多 GPU 运算
  1. 需要进一步提升计算速度。
  2. 需要计算更大的任务，内存需求超出单个 GPU
  3. 达到更高的能耗效率。
# 多 GPU 按照计算模式可以分为哪几类
  1. 单进程
    1. 单线程 + Single GPU
    2. 多线程 + Single GPU
  2. 多进程
    1. +Single GPU
    2. +Multiple GPU
# 简述统一虚拟地址 UVA 及其作用
UVA 为系统中的所有处理器（如 CPU 和 GPU）提供了一个统一的、连续的虚拟地址空间。这意味着无论从 CPU 还是 GPU 的角度来看，它们 ++ 都可以使用相同的虚拟地址 ++ 来访问系统内存和特定于设备的内存（如 GPU 的显存）
简化编程，方便内存管理。
GPU 可以定位另一个 GPU 的缓存位置。
# 简述 OpenCL 与 CUDA 的区别是什么
一、开发厂商与通用性
  1. OpenCL：
    • 由 苹果公司 发起，目前由 Khronos Group 维护。
    • 是一个开放的、跨平台的标准，支持多种硬件厂商的设备，包括 AMD、NVIDIA、Intel 的显卡以及各种移动设备的处理器等，具有广泛的硬件兼容性。
  2. CUDA：
    • 由 NVIDIA 公司开发。
    • 主要针对 NVIDIA 的 GPU，对自家硬件的优化程度高，但在其他硬件平台上无法直接使用。
二、编程模型
  1. OpenCL：
    • 基于任务并行和数据并行的混合模型，相对较为复杂。
    • 程序员需要明确地管理设备、内存对象、命令队列等底层概念，具有较高的灵活性，但也增加了编程的难度。
  2. CUDA：
    • 以数据并行为主，编程模型相对简单直观。
    • 提供了高级的编程语言扩展（如 C++ 扩展），程序员可以更方便地进行 GPU 编程，开发效率相对较高。
三、语言支持
  1. OpenCL：
    • 支持多种编程语言，如 C 、 C++ 、 Java 等，可以通过不同的语言绑定来使用 OpenCL。
  2. CUDA：
    • 主要使用 C/C++ 和一些特定的 CUDA 扩展。虽然也有其他语言的绑定，但相对较少。
四、生态系统与社区支持
  1. OpenCL：
    • 由于其跨平台性，拥有广泛的硬件支持和较大的社区，但在某些特定领域的生态系统可能不如 CUDA 完善。
  2. CUDA：
    • NVIDIA 对 CUDA 投入了大量资源进行开发和优化，拥有丰富的库和工具，如 cuDNN（用于深度学习）等，在深度学习、科学计算等领域有强大的生态系统和广泛的社区支持。
# OpenCL 有哪几种模型构成
平台模型 、执行模型 、内存模型 、编程模型
# 简述 OpenCL 编程的基本步骤都有哪些
  1. Selecting a Platform
  2. Contexts
    • A context refers to the environment for managing OpenCLobjects and resources
  3. Command Queues
  4. Memory Objects
  5. Creating buffers
  6. Transferring Data
# OpenCL 程序调试都有哪些工具
GDB（CPU_COMPILER_OPTIONS=“-g”）
GPU Printf
AMD GPUs support printing during execution using printf() NVIDIA does not currently support printing for OpenCL kernels (though they do with CUDA/C)
gDEBugger
# 简述 OpenCL 性能优化常见的方法
线性映射
  • 计算任务分配到不同计算单元
  • 使线程的命中率尽可能高
设备占用
寄存器，本地内存、线程资源
充分利用计算设备资源，以提高程序的性能和效率
合理地分配内存、选择合适的设备类型（如 GPU、CPU 或其他加速器）
向量化
在 OpenCL 中，可以使用向量数据类型和相应的指令来同时处理多个数据元素，从而提高计算效率。
# CUDA 编程模型
/**
 * Hello world!
*/
//int main(void){
//    printf("Hello World!\n");
//    return 0;
//}
/***************/
__global__ void mykernel(void){
}
int main(void){
    mykernel<<<1,1>>>();
    printf("Hello World!\n");
    return 0;
}
# GPU 求和计算
__global__ void add(int *a, int *b, int *c){
    	*c = *a + *b;
}
//a, b, c 指向 device 内存
// 需要在 GPU 上分配内存空间
int main(void){
    int a, b, c;
    int *d_a, *d_b, *d_c;
    int size = sizeof(int);
    
    cudaMalloc((void **)&d_a, device);
    cudaMalloc((void **)&d_b, device);
    cudaMalloc((void **)&d_c, device);
    
    a = 2;
    b = 7;
    
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
    
    add<<<1,1>>>(d_a, d_b, d_c);
    
    cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
    
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}
__global__ void add(int *a, int *b, int *c) {
 	c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}
#define N 512
int main(void){
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = N * sizeof(int);
    
    a = (int *)malloc(size); random_ints(a, N);
    b = (int *)malloc(size); random_ints(b, N);
    c = (int *)malloc(size);
    
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);
    
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
    
    add<<<N,1>>>(d_a, d_b, d_c); // N blocks, 1 thread per block
    
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
    
    free(a); free(b); free(c);
 	cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}
__global__ void add(int *a, int *b, int *c){
    c[threadIdx.x]=a[threadIdx.x] + b[threadIdx.x];
}
#define N 512
int main(void){
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = N * sizeof(int);
    
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);
    
    a = (int *)malloc(size); random_ints(a, N);
 	b = (int *)malloc(size); random_ints(b, N);
 	c = (int *)malloc(size);
    
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
    add<<<1,N>>>(d_a, d_b, d_c);
    
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
    free(a); free(b); free(c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}
__global__ void add(int *a, int *b, int *c){
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    c[index] = a[index] + b[index];
}
#define N (2048*2048)
#define THREADS_PER_BLOCK 512
int main(void){
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = N*sizeof(int);
    
    cudaMalloc((void **)&d_a, size);
 	cudaMalloc((void **)&d_b, size);
 	cudaMalloc((void **)&d_c, size);
    
    a = (int *)malloc(size); random_ints(a, N);
 	b = (int *)malloc(size); random_ints(b, N);
 	c = (int *)malloc(size);
    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
 	cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
    
    add<<<N/THREADS_PER_BLOCK，THREADS_PER_BLOCK>>>(d_a, d_b, d_c);
    
    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
    free(a); free(b); free(c);
 	cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
 	return 0;
}
__global__ void add(int *a, int *b, int *c, int n){
    int index = threadIdx.x + blockIdx.x*blockDim.x;
    if(index < n){
        c[index] = a[index] + c[index];
    }
}
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
 {
 int i = threadIdx.x;
 int j = threadIdx.y;
 C[i][j] = A[i][j] + B[i][j];
 }
 int main()
 {
 ...
 // Kernel invocation with one block of N * N * 1 threads
 int numBlocks = 1;
 dim3 threadsPerBlock(N, N);
 MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
 ...
 }
__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N])
 {
 int i = blockIdx.x * blockDim.x + threadIdx.x;
 int j = blockIdx.y * blockDim.y + threadIdx.y;
 if (i < N && j < N)
 C[i][j] = A[i][j] + B[i][j];
 }
 int main()
 {
 ...
 // 先定义每个 block 中线程 “形状” ，再计算 block 的 “形状”
 dim3 threadsPerBlock(16, 16);
 dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
 MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
 ...
 }
# OpenCL 编程模型
# 实例 —— 向量加法
__kernel void VectorAdd(__global const float* a,
                        __global const float* b,
                       	__global float* c,
                       	int iNumElements)
{
    int iGid = get_global_id(0);
    
    if(iGid >= iNumElements){
        return;
    }
    
    c[iGid] = a[iGid] + b[iGid];
}
int main(int argc, char** argv){
    localWorkSize = 256;
    globalWorkSize = RoundUp(localWorkSize, iNumElements);
    
    srcA = malloc(sizeof(cl_float) * globalWorkSize);
    srcB = malloc(sizeof(cl_float) * globalWorkSize);
    dst = malloc(sizeof(cl_float) * globalWorkSize);
    
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(0, 1, &device, NULL, NULL, &ciErr1);
    
    cqCommandQueue = clCreateCommandQueue(context, device, 0, &ciErr1);
    
    dsrcA = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float) * globalWorkSize, ...);
    dsrcB = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float) * globalWorkSize, ...);
    dDst = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float) * globalWorkSize, ...);
    
    clEnqueueWriteBuffer(cqCommandQueue, dsrcA, CL_FALSE, 0, sizeof(cl_float) * globalWorkSize, srcA, ...);
    clEnqueueWriteBuffer(cqCommandQueue, dsrcB, CL_FALSE, 0, sizeof(cl_float) * globalWorkSize, srcB, ...);
    
    progSrc = oclLoadProgSource("VectorAdd.cl", ...);
    prog = clCreateProgramWithSource(context, 1, progSrc, ...);
    clBuildProgram(prog, 0, NULL, NULL, NULL, NULL);
    
    kernel = clCreateKernel(prog, "VectorAdd", &ciErr1);
    
    clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&dSrcA);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&dSrcB);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&dDst);
    clSetKernelArg(kernel, 3, sizeof(cl_int), (void*)&iNumElements);
    
    clEnqueueNDRangeKernel(cqCommandQueue, kernel, 1, NULL, &globalWorkSize, &localWorkSize);
    
    clEnqueueReadBuffer(cqCommandQueue, dDst, ...);
    Cleanup();
    return 0;
    
}