c - 使用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，我也不确定这是否会超出扭曲的大小，并且肯定不会超出线程 block 的大小，但希望您已经意识到这一点(如果没有，请阅读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属于这里。