使用 cuda 编程的奇偶排序

Question

我正在尝试用 cuda-c 语言实现奇偶排序程序。但是，每当我将 0 作为输入数组中的元素之一时，结果数组没有正确排序。但是，在其他情况下，它适用于其他输入。我不明白代码有什么问题.这是我的代码：

#include<stdio.h>
#include<cuda.h>
#define N 5

__global__ void sort(int *c,int *count)
{
    int l;
    if(*count%2==0)
          l=*count/2;
    else
         l=(*count/2)+1;
    for(int i=0;i<l;i++)
    {
            if(threadIdx.x%2==0)  //even phase
            {
                if(c[threadIdx.x]>c[threadIdx.x+1])
                {
                    int temp=c[threadIdx.x];
                    c[threadIdx.x]=c[threadIdx.x+1];
                    c[threadIdx.x+1]=temp;
                }

            __syncthreads();
            }
            else     //odd phase
            {
                if(c[threadIdx.x]>c[threadIdx.x+1])
                {
                    int temp=c[threadIdx.x];
                    c[threadIdx.x]=c[threadIdx.x+1];
                    c[threadIdx.x+1]=temp;
                }

            __syncthreads();
            }
    }//for

}



int main()
{int a[N],b[N],n;
    printf("enter size of array");
    scanf("%d",&n);
    print("enter the elements of array");
  for(int i=0;i<n;i++)
  {
    scanf("%d",&a[i]);
  }
  printf("ORIGINAL ARRAY : \n");
  for(int i=0;i<n;i++)
          {

          printf("%d ",a[i]);
          }
  int *c,*count;
  cudaMalloc((void**)&c,sizeof(int)*N);
  cudaMalloc((void**)&count,sizeof(int));
  cudaMemcpy(c,&a,sizeof(int)*N,cudaMemcpyHostToDevice);
  cudaMemcpy(count,&n,sizeof(int),cudaMemcpyHostToDevice);
  sort<<< 1,n >>>(c,count);
  cudaMemcpy(&b,c,sizeof(int)*N,cudaMemcpyDeviceToHost);
  printf("\nSORTED ARRAY : \n");
  for(int i=1;i<=n;i++)
      {
         printf("%d ",b[i]);
      }

}

Answer 1

您的内核代码有两个我可以看到的主要错误：

1. 在奇数阶段（对于偶数长度数组，或者对于奇数长度数组的偶数阶段），您的最后一个线程将在c[threadIdx.x+1] 处索引超出范围。例如，对于 4 个线程，它们的编号为 0、1、2、3。线程 3 很奇怪，但是如果您访问 c[3+1]，那不是您的数组中定义的元素。我们可以通过限制每个阶段在除最后一个线程之外的所有线程上工作来解决此问题。

2. 您在不允许所有线程到达屏障的条件语句中使用__syncthreads()。这是一个编码错误。阅读documentation。我们可以通过调整条件区域内的代码来解决这个问题。

在主代码中，您的最终打印输出语句索引不正确：

for(int i=1;i<=n;i++)

应该是：

for(int i=0;i<n;i++)

这里还有错别字：

print("enter the elements of array");

我认为应该是printf。

以下代码已修复上述错误，并且对于长度为 5 的数组（您对 N 的硬编码限制）来说似乎可以正确运行。即使您增加了N，我也不确定这是否会超出经纱的大小，当然也不会超出线程块的大小，但希望您已经意识到这一点（如果没有，请阅读doc link about __syncthreads()） .

“固定”代码：

#include<stdio.h>
#include<cuda.h>
#define N 5

#define intswap(A,B) {int temp=A;A=B;B=temp;}

__global__ void sort(int *c,int *count)
{
    int l;
    if(*count%2==0)
          l=*count/2;
    else
         l=(*count/2)+1;
    for(int i=0;i<l;i++)
    {
            if((!(threadIdx.x&1)) && (threadIdx.x<(*count-1)))  //even phase
            {
                if(c[threadIdx.x]>c[threadIdx.x+1])
                  intswap(c[threadIdx.x], c[threadIdx.x+1]);
            }

            __syncthreads();
            if((threadIdx.x&1) && (threadIdx.x<(*count-1)))     //odd phase
            {
                if(c[threadIdx.x]>c[threadIdx.x+1])
                  intswap(c[threadIdx.x], c[threadIdx.x+1]);
            }
            __syncthreads();
    }//for

}



int main()
{int a[N],b[N],n;
    printf("enter size of array");
    scanf("%d",&n);
    if (n > N) {printf("too large!\n"); return 1;}
    printf("enter the elements of array");
  for(int i=0;i<n;i++)
  {
    scanf("%d",&a[i]);
  }
  printf("ORIGINAL ARRAY : \n");
  for(int i=0;i<n;i++)
          {

          printf("%d ",a[i]);
          }
  int *c,*count;
  cudaMalloc((void**)&c,sizeof(int)*N);
  cudaMalloc((void**)&count,sizeof(int));
  cudaMemcpy(c,&a,sizeof(int)*N,cudaMemcpyHostToDevice);
  cudaMemcpy(count,&n,sizeof(int),cudaMemcpyHostToDevice);
  sort<<< 1,n >>>(c,count);
  cudaMemcpy(&b,c,sizeof(int)*N,cudaMemcpyDeviceToHost);
  printf("\nSORTED ARRAY : \n");
  for(int i=0;i<n;i++)
      {
         printf("%d ",b[i]);
      }

  printf("\n");
}

关于proper cuda error checking 的常规独奏属于这里。