将通过设计一个示例来开始我们的GPU并行性探索，该示例与大多数可用GPGPU文档中的示例有所不同，该示例主要由具有较强图形或科学背景的人编写。 这些示例中的大多数都谈论矢量加法或其他一些数学构造。 让我们来研究一个有点类似于业务情况的示例。 因此，让我们首先假设我们有一个产品列表，每个产品都有一个代码标签和一个价格，并且我们想对所有代码为“ abc”的产品施加10％的开销。
我们将首先在C语言中实现示例，以在CPU和GPU处理之间进行一些性能评估。 之后，我们当然会在Java中实现相同的功能，但是我们将避免进行任何测量，因为它们在Java中有些棘手，因为我们必须考虑垃圾收集，及时编译等问题。
//============================================================================
// Name : StoreDiscountExample.cu
// Author : Spyros Sakellariou
// Version : 1.0
// Description : The Good the Bad and the Ugly
//============================================================================
#include <iostream>
#include <sys/time.h>
typedef struct {
char code[3];
float listPrice;
} product;
void printProducts(long size, product * myProduct){
printf("Price of First item=%f,%s\n",myProduct[0].listPrice,myProduct[0].code);
printf("Price of Second item=%f,%s\n",myProduct[1].listPrice,myProduct[1].code);
printf("Price of Middle item=%f,%s\n",myProduct[(size-1)/2].listPrice,myProduct[(size-1)/2].code);
printf("Price of Almost Last item=%f,%s\n",myProduct[size-2].listPrice,myProduct[size-2].code);
printf("Price of Last item=%f,%s\n",myProduct[size-1].listPrice,myProduct[size-1].code);
}
float calculateMiliseconds (timeval t1,timeval t2) {
float elapsedTime;
elapsedTime = (t2.tv_sec - t1.tv_sec) * 1000.0;
elapsedTime += (t2.tv_usec - t1.tv_usec) / 1000.0;
return elapsedTime;
}
__global__ void kernel(long size, product *products)
{
long kernelid = threadIdx.x + blockIdx.x * blockDim.x;
while(kernelid < size) {
if (products[kernelid].code[0]=='a' && products[kernelid].code[1]=='b' && products[kernelid].code[2]=='c')
products[kernelid].listPrice*=1.10;
kernelid += blockDim.x * gridDim.x;
}
}
int main( int argc, char** argv)
{
timeval t1,t2;
cudaEvent_t eStart,eStop;
float elapsedTime;
long threads = 256;
long blocks = 1024;
long size = 9000000;
char *product1 = "abc";
char *product2 = "bcd";
product *myProduct;
product *dev_Product;
printf("blocks=%d x threads=%d total threads=%d total number of products=%d\n\n",blocks,threads,threads*blocks,size);
myProduct=(product*)malloc(sizeof(myProduct)*size);
cudaMalloc((void**)&dev_Product,sizeof(dev_Product)*size);
cudaEventCreate(&eStart);
cudaEventCreate(&eStop);
gettimeofday(&t1, NULL);
for (long i = 0; i<size; i++){
if (i%2==0)
strcpy(myProduct[i].code,product1);
else
strcpy(myProduct[i].code,product2);
myProduct[i].listPrice = i+1;
}
gettimeofday(&t2, NULL);
printf ( "Initialization time %4.2f ms\n", calculateMiliseconds(t1,t2) );
printProducts(size,myProduct);
cudaMemcpy(dev_Product,myProduct,sizeof(dev_Product)*size,cudaMemcpyHostToDevice);
cudaEventRecord(eStart,0);
kernel<<<blocks,threads>>>(size,dev_Product);
cudaEventRecord(eStop,0);
cudaEventSynchronize(eStop);
cudaMemcpy(myProduct,dev_Product,sizeof(dev_Product)*size,cudaMemcpyDeviceToHost);
cudaEventElapsedTime(&elapsedTime,eStart,eStop);
printf ( "\nCuda Kernel Time=%4.2f ms\n", elapsedTime );
printProducts(size,myProduct);
gettimeofday(&t1, NULL);
long j=0;
while (j < size){
if (myProduct[j].code[0]=='a' && myProduct[j].code[1]=='b' && myProduct[j].code[2]=='c')
myProduct[j].listPrice*=0.5;
j++;
}
gettimeofday(&t2, NULL);
printf ( "\nCPU Time=%4.2f ms\n", calculateMiliseconds(t1,t2) );
printProducts(size,myProduct);
cudaFree(dev_Product);
free(myProduct);
}在第11至14行中，定义了一个结构，其中包含我们的产品，该产品带有用于产品代码的字符数组和用于其价格的浮点数。
在第16和24行中，定义了两种实用程序方法，一种打印某些产品（以便我们查看是否完成工作），另一种将原始日期差转换为毫秒。 请注意，使用标准C时钟功能将无法使用，因为其粒度不足以测量毫秒。
第33行是编写我们的内核的位置。 与上一篇文章相比，这看起来有些复杂，因此让我们进一步剖析…
在第35行中，我们定义了kernelid参数。 此参数将保存正在执行的线程的唯一线程ID。 CUDA为每个线程分配一个线程ID和块ID号，该ID仅对自己的维度唯一。 在我们的示例中，我们指示GPU启动256个线程和1024个块，因此实际上GPU将执行262144个线程。 很好！ 尽管CUDA在执行过程中为我们提供了threadIdx.x和blockIdx.x参数，但是我们需要手动创建唯一的线程ID才能知道我们当前所在的线程。唯一的线程ID需要从0到262143开始，因此，我们可以通过将要执行的每个块的线程数（使用CUDA参数blockDim.x ） 乘以当前块并将其添加到当前线程中来轻松创建它，从而：
唯一线程ID =当前线程ID +当前块ID *每个块的线程数
如果您想继续阅读，您已经意识到，尽管26.2万个线程令人印象深刻，但我们的数据集由900万个项目构成，因此我们的线程一次需要处理多个项目。 为此，我们在第36行中设置了一个循环，以检查我们的线程ID是否不会超出我们的产品数据数组。 循环使用线程id作为其索引，但是我们使用以下公式对其进行递增：
索引增量+ =每块线程数*块总数
因此，每个线程将执行每个循环9百万/ 262千次，这意味着它将处理约34个项目。
其余的内核代码非常简单且易于说明，每当我们找到带有代码“ abc”的乘积时，便将其乘以1.1（我们的10％开销）。 注意，不能在内核内部使用strcpy函数。 如果尝试将出现编译错误。 不太好！
在第45和46行进入主要功能时，我们定义了两个C计时器（ t1和t2 ）和两个CUDA事件计时器（ eStart和eStop ）。 我们需要CUDA计时器，因为内核是异步执行的，并且我们的内核函数会立即返回，而我们的计时器将仅测量完成函数调用所花费的时间。 内核代码立即返回的事实意味着我们允许CPU在GPU代码执行期间执行其他任务。 这也很好！
计时器后面的参数不言自明：我们定义了线程数，块数，产品数组的大小等。myproduct指针将用于CPU处理，而dev_product指针将用于GPU处理。
在第58至61行中，我们为myproduct和dev_product分配了RAM和GPU内存，并且我们还创建了CUDA计时器来帮助我们测量内核执行时间。
在第63至73行中，我们使用代码和价格初始化myproduct ，然后打印CPU完成任务所花费的时间。 我们还从阵列中打印一些样品产品，以确保正确完成作业。
在第74到85行中，我们将产品数组复制到GPU内存，执行内核以指示要执行的线程和块的数量，然后将结果复制回myproduct数组。 我们打印执行内核和一些示例产品所花费的时间，以确保我们再次正确完成了工作。
最后，在第88至99行中，我们让CPU执行与GPU相似的过程，即对GPU增加了开销的所有产品进行50％的折扣。 我们打印执行CPU任务所花费的时间，并打印一些示例产品以确保完成工作。
让我们编译并运行以下代码：
# nvcc StoreDiscountExample.cu -o StoreDiscountExample
# ./StoreDiscountExample
blocks=1024 x threads=256 total threads=262144 total number of products=9000000
Initialization time 105.81 ms
Price of First item=1.000000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=8999999.000000,abc
Price of Last item=9000000.000000,bcd
Cuda Kernel Time=1.38 ms
Price of First item=1.100000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=9899999.000000,abc
Price of Last item=9000000.000000,bcd
CPU Time=59.58 ms
Price of First item=0.550000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=4949999.500000,abc
Price of Last item=9000000.000000,bcd
#哇！ GPU花费了1.38毫秒来完成CPU花费了59.58毫秒（当然，具体数字取决于您的硬件）。 很好！
稍等片刻！ 在决定删除所有代码并开始重新编写CUDA中的所有内容之前，有一个陷阱：我们忽略了一些严重的问题，那就是要衡量将900万条记录从RAM复制到GPU内存再返回所需的时间。 这是从第74行到第85行的代码，更改为使用计时器来衡量将产品列表复制到GPU内存以及从GPU内存复制产品列表的时间：
gettimeofday(&t1, NULL);
cudaMemcpy(dev_Product,myProduct,sizeof(dev_Product)*size,cudaMemcpyHostToDevice);
cudaEventRecord(eStart,0);
kernel<<<blocks,threads>>>(size,dev_Product);
cudaEventRecord(eStop,0);
cudaEventSynchronize(eStop);
cudaMemcpy(myProduct,dev_Product,sizeof(dev_Product)*size,cudaMemcpyDeviceToHost);
gettimeofday(&t2, NULL);
printf ( "\nCuda Total Time=%4.2f ms\n", calculateMiliseconds(t1,t2));
cudaEventElapsedTime(&elapsedTime,eStart,eStop);
printf ( "Cuda Kernel Time=%4.2f ms\n", elapsedTime );
printProducts(size,myProduct);让我们编译并运行以下代码：
# nvcc StoreDiscountExample.cu -o StoreDiscountExample
# ./StoreDiscountExample
blocks=1024 x threads=256 total threads=262144 total number of products=9000000
Initialization time 108.31 ms
Price of First item=1.000000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=8999999.000000,abc
Price of Last item=9000000.000000,bcd
Cuda Total Time=55.13 ms
Cuda Kernel Time=1.38 ms
Price of First item=1.100000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=9899999.000000,abc
Price of Last item=9000000.000000,bcd
CPU Time=59.03 ms
Price of First item=0.550000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=4949999.500000,abc
Price of Last item=9000000.000000,bcd
#请注意，CUDA总时间为55毫秒，仅比在单个线程中使用CPU快4毫秒。 这是不好的！
因此，尽管GPU在并行任务执行方面非常快，但是当我们在RAM和GPU内存之间来回复制项目时，代价是很大的。 有一些高级技巧，例如可以使用直接内存访问，但道德的道理是在决定使用GPU时必须非常小心。 如果您的算法需要大量数据移动，那么GPGPU可能不是答案。
由于我们已经完成了性能测试，因此让我们看一下如何使用jcuda实现相同的功能。
这是Java部分的代码：
import static jcuda.driver.JCudaDriver.*;
import jcuda.*;
import jcuda.driver.*;
import jcuda.runtime.JCuda;
public class StoreDiscountExample {
public static void main(String[] args) {
int threads = 256;
int blocks = 1024;
final int size = 9000000;
byte product1[] = "abc".getBytes();
byte product2[] = "bcd".getBytes();
byte productList[] = new byte[size*3];
float productPrices[] = new float[size];
long size_array[] = {size};
cuInit(0);
CUcontext pctx = new CUcontext();
CUdevice dev = new CUdevice();
cuDeviceGet(dev, 0);
cuCtxCreate(pctx, 0, dev);
CUmodule module = new CUmodule();
cuModuleLoad(module, "StoreDiscountKernel.ptx");
CUfunction function = new CUfunction();
cuModuleGetFunction(function, module, "kernel");
int j=0;
for (int i = 0; i<size; i++){
j=i*3;
if (i%2==0) {
productList[j]=product1[0];
productList[j+1]=product1[1];
productList[j+2]=product1[2];
}
else {
productList[j]=product2[0];
productList[j+1]=product2[1];
productList[j+2]=product2[2];
}
productPrices[i] = i+1;
}
printSamples(size, productList, productPrices);
CUdeviceptr size_dev = new CUdeviceptr();
cuMemAlloc(size_dev, Sizeof.LONG);
cuMemcpyHtoD(size_dev, Pointer.to(size_array), Sizeof.LONG);
CUdeviceptr productList_dev = new CUdeviceptr();
cuMemAlloc(productList_dev, Sizeof.BYTE*3*size);
cuMemcpyHtoD(productList_dev, Pointer.to(productList), Sizeof.BYTE*3*size);
CUdeviceptr productPrice_dev = new CUdeviceptr();
cuMemAlloc(productPrice_dev, Sizeof.FLOAT*size);
cuMemcpyHtoD(productPrice_dev, Pointer.to(productPrices), Sizeof.FLOAT*size);
Pointer kernelParameters = Pointer.to(
Pointer.to(size_dev),
Pointer.to(productList_dev),
Pointer.to(productPrice_dev)
);
cuLaunchKernel(function,
blocks, 1, 1,
threads, 1, 1,
0, null,
kernelParameters, null);
cuMemcpyDtoH(Pointer.to(productPrices), productPrice_dev, Sizeof.FLOAT*size);
printSamples(size, productList, productPrices);
JCuda.cudaFree(productList_dev);
JCuda.cudaFree(productPrice_dev);
JCuda.cudaFree(size_dev);
}
public static void printSamples(int size, byte[] productList, float[] productPrices) {
System.out.print(String.copyValueOf(new String(productList).toCharArray(), 0, 3));System.out.println(" "+productPrices[0]);
System.out.print(String.copyValueOf(new String(productList).toCharArray(), 3, 3));System.out.println(" "+productPrices[1]);
System.out.print(String.copyValueOf(new String(productList).toCharArray(), 6, 3));System.out.println(" "+productPrices[2]);
System.out.print(String.copyValueOf(new String(productList).toCharArray(), 9, 3));System.out.println(" "+productPrices[3]);
System.out.print(String.copyValueOf(new String(productList).toCharArray(), (size-2)*3, 3));System.out.println(" "+productPrices[size-2]);
System.out.print(String.copyValueOf(new String(productList).toCharArray(), (size-1)*3, 3));System.out.println(" "+productPrices[size-1]);
}
} 从第14、15和16行开始，我们看到我们不能再使用产品结构或类。 实际上，在jcuda中，将作为参数传递给内核的所有内容都必须是一维数组。 因此，在第14行中，我们创建了表示为一维数组的字节的二维数组。 产品列表数组大小等于产品数量乘以每个产品代码的字节大小（在我们的情况下只有3个字节）。 我们还创建了另一个数组，以浮动价格存储产品价格，最后，我们的产品列表的大小也需要放入一维数组中。
我想到现在您可能已经猜到了我要说的是：这简直太丑了！
在第29至45行中，我们填充产品列表和产品价格数组，然后在调用内核函数之前，通过创建CUDA设备指针，分配内存并将数据复制到GPU内存，将它们传递给内核进行处理。
由于我们必须将所有内容都转换为基元的一维数组，因此我们的内核代码也需要进行一些更改：
extern "C"
__global__ void kernel(long *size, char *productCodes, float *productPrices)
{
long kernelid = threadIdx.x + blockIdx.x * blockDim.x;
long charIndex = kernelid*3;
while(kernelid < size[0]) {
if (productCodes[charIndex]=='a' && productCodes[charIndex+1]=='b' && productCodes[charIndex+2]=='c')
productPrices[kernelid]*=1.10;
kernelid += blockDim.x * gridDim.x;
charIndex = kernelid*3;
}
}唯一的区别是，我们将kernelid索引乘以3，以便在productCodes数组中找到正确的起始字符。
让我们编译并运行Java示例：
# nvcc -ptx StoreDiscountKernel.cu -o StoreDiscountKernel.ptx
# javac -cp ~/GPGPU/jcuda/JCuda-All-0.4.0-beta1-bin-linux-x86_64/jcuda-0.4.0-beta1.jar StoreDiscountExample.java
#java -cp ~/GPGPU/jcuda/JCuda-All-0.4.0-beta1-bin-linux-x86_64/jcuda-0.4.0-beta1.jar StoreDiscountExample
abc 1.0
bcd 2.0
abc 3.0
bcd 4.0
abc 8999999.0
bcd 9000000.0
abc 1.1
bcd 2.0
abc 3.3
bcd 4.0
abc 9899999.0
bcd 9000000.0
#尽管代码很丑陋，但其工作方式与C语言类似。
因此，这里总结了我们在GPU处理和jcuda方面的经验：
良好 ：非常快的性能（我的硬件是AMD Phenom II四核3.4Ghz CPU和具有336核的NVIDIA Geforce GTX 560）
良好 ：异步操作让CPU执行其他任务
坏 ：内存副本会带来可观的性能损失
丑陋 ：如果您想从Java内部执行CUDA内核，Jcuda无疑是有用的，但是必须将所有内容都转换为基元的一维数组实在不方便。
在我们之前的文章中，有一些关于OpenCL（GPGPU的CUDA替代品）的Java工具的非常有趣的评论。 在下一篇文章中，我们将研究这些工具，看看它们是否比Jcuda更“漂亮”。
参考： GPGPU与我们的W4G合作伙伴 Spyros Sakellariou的 Jcuda the Good，Bad和…Ugly 一起使用 。
相关文章 ：
翻译自: https://www.javacodegeeks.com/2011/10/gpgpu-with-jcuda-good-bad-and-ugly.html
扫一扫
423
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
