1

我在 Visual Studio 2010 Professional 环境中的 Windows 7 64 位上运行 CUDA 4.2

首先，我运行以下代码：

// include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std; 

//kernel function
__global__ 
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid==0)
       d_bStopPtr[tid]=0;
    else if(tid<count)
    {
       d_bPtr[tid]=tid;
// only if the arrary cell before it is 0, then change it to 0 too
        if (d_bStopPtr[tid-1]==0 )
           d_bStopPtr[tid]=0;

    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

    int threadsPerBlock = prop.maxThreadsDim[0];
    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
    while(d_bStop[count-1])
    {
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
    }
//copy device back to host again
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

但是，每次我都需要检查while条件，但速度很慢。所以我正在考虑检查内核中这个设备向量的条件，并像这样更改代码：

// include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std; 

//kernel function
__global__ 
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid==0)
    d_bStopPtr[tid]=0;
else if(tid<count)
    {
// if the last cell of the arrary is still not 0 yet, repeat
        while(d_bStopPtr[count-1])
        {
            d_bPtr[tid]=tid;
// only if the arrary cell before it is 0, then change it to 0 too
            if (d_bStopPtr[tid-1]==0 )
                d_bStopPtr[tid]=0;
        }
    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

    int threadsPerBlock = prop.maxThreadsDim[0];
    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
//copy device back to host again
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

但是，第二个版本总是会导致显卡和电脑死机。你能帮我加快第一个版本的速度吗？如何检查内核内部的条件然后跳出并停止内核？

4

1 回答 1

2

您基本上是在寻找全局线程同步行为。这是 GPU 编程中的禁忌。理想情况下，每个线程块都是独立的，可以根据自己的数据和处理完成工作。创建依赖于其他线程块的结果来完成其工作的线程块会产生死锁情况的可能性。假设我有一个带有 14 个 SM（线程块执行单元）的 GPU，并假设我创建了 100 个线程块。现在假设线程块 0-13 正在等待线程块 99 释放锁（例如，将零值写入特定位置）。现在假设前 14 个线程块开始在 14 个 SM 上执行，可能是循环，在锁定值上旋转。GPU 中没有任何机制可以保证线程块 99 会首先执行，甚至根本不会执行，

让我们不要讨论“强制驱逐线程块 0-13 的 GMEM 停顿会怎样”的问题，因为这些都不能保证线程块 99 将在任何时候获得优先执行。唯一能保证线程块 99 将执行的是其他线程块的耗尽（即完成）。但是如果其他线程块正在旋转，等待线程块 99 的结果，那可能永远不会发生。

良好的前向兼容、可扩展的 GPU 代码依赖于独立的并行工作。因此，建议您重新设计算法以使您尝试完成的工作独立，至少在线程块间级别。

如果您必须进行全局线程同步，内核启动是唯一真正保证的点，因此您的第一种方法是工作方法。

为了解决这个问题，研究如何在 GPU 上实现缩减算法可能会很有用。各种类型的归约对所有线程都有依赖性，但通过创建中间结果，我们可以将工作分解为独立的部分。然后可以使用多内核方法（或一些其他更高级的方法）聚合独立的部分，以加快相当于串行算法的速度。

您的内核实际上并没有做太多事情。它设置一个数组等于它的索引，即 a[i] = i; 并且它将另一个数组设置为全零（尽管顺序）b [i] = 0;。要显示您的第一个代码“加速”的示例，您可以执行以下操作：

    // include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;

//kernel function
__global__
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid<count)
    {
      d_bPtr[tid]=tid;
      while(d_bStopPtr[tid]!=0)
// only if the arrary cell before it is 0, then change it to 0 too
        if (tid==0) d_bStopPtr[tid] =0;
        else if (d_bStopPtr[tid-1]==0 )
               d_bStopPtr[tid]=0;
      tid += blockDim.x;
    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

//    int threadsPerBlock = prop.maxThreadsDim[0];
    int threadsPerBlock = 32;
//    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
    int blocksPerGrid = 1;
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
//    while(d_bStop[count-1])
//    {
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
//    }
//copy device back to host again
    cudaDeviceSynchronize();
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

在我的机器上，这将执行时间从 10 秒加快到几乎是瞬时的（远小于 1 秒）。请注意，这不是 CUDA 编程的一个很好的例子，因为我只启动一个 32 个线程的块。这还不足以有效地利用机器。但是你的内核所做的工作是如此微不足道，以至于我不确定一个好的例子会是什么。我可以创建一个内核，将一个数组设置为其索引 a[i]=i; 另一个数组为零 b[i]=0; 全部并行。那会更快，我们可以这样使用整台机器。

于 2012-10-16T05:19:32.813 回答