1

我将 CUDA 用于迭代 Karatsuba 算法，我想问一下，为什么计算出来的一条线总是不同的。

首先，我实现了这个函数，它总是正确地计算结果：

__global__ void kernel_res_main(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;


        TYPE end = (i + 1) / 2;


        for(TYPE inner = start; inner < end; inner++){
            result[i] += ( A[inner] + A[i - inner] ) * ( B[inner] + B[i - inner] );
            result[i] -= ( D[inner] + D[i-inner] );
        }
    }
}

现在我想使用 2D 网格并使用 CUDA 进行 for 循环，所以我将函数更改为：

__global__ void kernel_res_nested(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    TYPE rtmp = result[i];

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;
        TYPE end = (i + 1) >> 1;

        if(j >= start && j <= end ){

           // WRONG 
           rtmp += ( A[j] + A[i - j] ) * ( B[j] + B[i - j] ) - ( D[j] + D[i - j] );
        }
    }

    result[i] = rtmp;
}

我这样调用这个函数：

dim3 block( 32, 8 );
dim3 grid( (resultSize+1/32) , (resultSize+7/8) );
kernel_res_nested <<<grid, block>>> (devA, devB, devD, devResult, size, resultSize);

结果总是错误的，总是不同的。我不明白为什么第二个实现是错误的并且总是计算错误的结果。我看不出有任何与数据依赖相关的逻辑问题。有谁知道我该如何解决这个问题？

4

1 回答 1

0

对于这样的问题，您应该提供一个 MCVE。（请参阅此处的第 1 项）例如，我不知道 表示什么类型TYPE，这对我将提出的解决方案的正确性很重要。

在您的第一个内核中，整个网格中只有一个线程在读取和写入 location result[i]。但是在您的第二个内核中，您现在有多个线程写入该result[i]位置。他们彼此冲突。CUDA 没有指定线程运行的顺序，有些线程可能在其他线程之前、之后或同时运行。在这种情况下，某些线程可能会与其他线程同时读取result[i]。然后，当线程写入它们的结果时，它们将不一致。它可能会因运行而异。你有一个竞争条件（执行顺序依赖，而不是数据依赖）。

解决这个问题的规范方法是使用归约技术。

但是为了简单起见，我建议原子可以帮助您解决问题。根据您所展示的内容，这更容易实现，并将有助于确认竞态条件。在那之后，如果你想尝试一种减少方法，有很多教程（上面链接了一个），cuda标签上有很多关于它的问题。

您可以将内核修改为类似这样，以理清竞争条件：

__global__ void kernel_res_nested(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;
        TYPE end = (i + 1) >> 1;

        if(j >= start && j < end ){ // see note below

           atomicAdd(result+i, (( A[j] + A[i - j] ) * ( B[j] + B[i - j] ) - ( D[j] + D[i - j] )));
        }
    }

}

请注意，根据您的 GPU 类型和TYPE您使用的实际类型，这可能无法按原样工作（可能无法编译）。但是由于您以前曾用作TYPE循环变量，因此我假设它是整数类型，并且atomicAdd应该可以使用这些必要的类型。

其他一些评论：

  1. 这可能不会为您提供您期望的网格大小：

    dim3 grid( (resultSize+1/32) , (resultSize+7/8) );

    我认为通常的计算是：

    dim3 grid( (resultSize+31)/32, (resultSize+7)/8 );

  2. 我总是推荐正确的 CUDA 错误检查和运行你的代码cuda-memcheck，任何时候你遇到 CUDA 代码的问题，以确保没有运行时错误。

  3. 在我看来它也是这样的：

    if(j >= start && j <= end ){

    应该是这样的：

    if(j >= start && j < end ){

    以匹配您的 for 循环范围。我也在做一个size小于resultSize（同样，MCVE 会有所帮助）的假设。

于 2018-04-25T20:44:43.850 回答