0

我想在CUDA C中并行化一个函数，它将计算所有向量的和等于向量元素和不大于k的元素。例如，如果向量元素的数量 n 为 5，sum=10 且 k=3，则满足此条件的向量数量为 101。我已经在 CUDA C 中制作了此函数，但问题是当块和线程大于 1。我知道问题出在for循环中，我应该更改它，但我不知道从哪里开始。当我调用块和线程等于一的函数时，该函数以经典方式工作，一切都很好，但在这种情况下，函数没有并行化。

该程序的源代码是：

//function that count number of vectors
__device__ void count(int *vector, int *total, int n, int s)
{
 int i,sum=0;
 for(i=blockIdx.x*blockDim.x+threadIdx.x;i<n;i+=blockDim.x*gridDim.x)
 {    
   sum+=vector[i];
   __syncthreads();
 }
 if(sum==s)
 {  
   total[0]=total[0]+1;
 }
}

//main function
__global__ void computeVectors(int *vector, int n, int kk, int s, int *total)
{
 int k=0;
 int j,i,next;

 while(1)
 {
  //this is the problem, in for cycle
  for(j=blockIdx.x*blockDim.x+threadIdx.x; j<=kk; j+=blockDim.x*gridDim.x)
  {
   vector[k]=j;
   count(vector, total, n, s);
   __syncthreads();
  }
  for(i=blockIdx.x*blockDim.x+threadIdx.x; i<n; i+=blockDim.x*gridDim.x)
  {

   if(vector[i]<kk)
      break;
  }  
  next=i;
  vector[next]++;
  for(i=blockIdx.x*blockDim.x+threadIdx.x; i<sledno; i+=blockDim.x*gridDim.x)
  {
   vector[i]=0;
   __syncthreads();
  }
  k=0;
  if(next>=n)
    break;
 }
}

int main()
{
  cudaError_t err = cudaSuccess;

  int n,k,sum;
  int counter=0;

  printf("Enter the length of vector n=");
  scanf("%d",&n);
  printf("Enter the max value of vector elements k=");
  scanf("%d",&k);
  printf("Enter the sum of vector elements sum=");
  scanf("%d",&sum);

  //initial vector with length n
  int *vec_h, *vec_d;
  size_t sizevec=n*sizeof(int);
  vec_h=(int *)malloc(sizevec);
  cudaMalloc((void **) &vec_d, sizevec);

  for(counter=0; counter<n; counter++)
  {
   vec_h[counter]=0;
  }
  cudaMemcpy(vec_d, vec_h, sizevec, cudaMemcpyHostToDevice);

  int *total_h, *total_d;
  size_t size=1*sizeof(int);
  total_h=(int *)malloc(size);
  cudaMalloc((void **) &total_d, size);
  total_h[0]=0;
  cudaMemcpy(total_d, total_h, size, cudaMemcpyHostToDevice);

  //calling the main function
  computeVectors<<<1, 1>>>(vec_d, n, k, sum, total_d);

  cudaThreadSynchronize(); 

  err = cudaGetLastError();
  if (err != cudaSuccess)
  {
    fprintf(stderr, "Error: %s!\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
  }
  cudaMemcpy(total_h, total_d, size, cudaMemcpyDeviceToHost);
  printf("Number of vectors that satisfy condition is %d\n", total_h[0]);

  free(vec_h); 
  cudaFree(vec_d);

  free(total_h); 
  cudaFree(total_d);

  return 0;
}
4

3 回答 3

1

这是一个示例蛮力程序，用于枚举所有可能的向量，然后测试每个向量的总和以查看它是否与所需的总和匹配。

  • 假设n=“数字”中的向量长度
  • 假设每个向量“数字”由一个unsigned数量表示
  • 假设k= 最大“数字”值 + 1
  • 向量空间的大小由k^给出n
  • 将此空间划分为要由每个线程处理的连续向量组： ( k^ n)/grid_size
  • 为每个线程生成起始向量（即每组中的起始向量）
  • 然后每个线程循环通过测试向量和并在必要时增加计数，然后“增加”向量，直到每个线程处理完分配给它的连续向量组

该程序：

#include <stdio.h>
#include <thrust/host_vector.h>
#include <sys/time.h>
#include <time.h>

#define MAX_N 12
#define nTPB 256
#define GRIDSIZE (32*nTPB)


#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


// thrust code is to quickly prototype a CPU based
// method for verification
int increment(thrust::host_vector<unsigned> &data, unsigned max){
  int pos = 0;
  int done = 0;
  int finished = 0;

  while(!done){
    data[pos]++;
    if (data[pos] >= max) {
      data[pos] = 0;
      pos++;
      if (pos >= data.size()){
        done = 1;
        finished = 1;
        }
      }
    else done = 1;
  }
  return finished;
}

__constant__ unsigned long powers[MAX_N];

__device__ unsigned vec_sum(unsigned *vector, int size){
  unsigned sum = 0;
  for (int i=0; i<size; i++) sum += vector[(i*nTPB)];
  return sum;
}

__device__ void create_vector(unsigned long index, unsigned *vector, int size){
  unsigned long residual = index;
  unsigned pos = size;
  while ((residual > 0) && (pos > 0)){
    unsigned long temp = residual/powers[pos-1];
    vector[(pos-1)*nTPB] = temp;
    residual -= temp*powers[pos-1];
    pos--;
    }
  while (pos>0) {
   vector[(pos-1)*nTPB] = 0;
   pos--;
   }
}
__device__ void increment_vector(unsigned *vector, int size, int k){
  int pos = 0;
  int done = 0;

  while(!done){
    vector[(pos*nTPB)]++;
    if (vector[pos*nTPB] >= k) {
      vector[pos*nTPB] = 0;
      pos++;
      if (pos >= size){
        done = 1;
        }
      }
    else done = 1;
  }
}

__global__ void find_vector_match(unsigned long long int *count, int k, int n, unsigned sum){
  __shared__ unsigned vecs[MAX_N *nTPB];
  unsigned *vec = &(vecs[threadIdx.x]);
  unsigned long idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < (k*powers[n-1])){
    unsigned long vec_count = 0;
    unsigned long vecs_per_thread = (k*powers[n-1])/(gridDim.x*blockDim.x);
    vecs_per_thread++;
    unsigned long vec_num = idx*vecs_per_thread;
    create_vector((vec_num), vec, n);
    while ((vec_count < vecs_per_thread) && (vec_num < (k*powers[n-1]))){
      if (vec_sum(vec, n) == sum) atomicAdd(count, 1UL);
      increment_vector(vec, n, k);
      vec_count++;
      vec_num++;
      }
   }
}

int main(){

// calculate on CPU first for verification
  struct timeval t1, t2, t3;
  int n, k, sum;
  printf("Enter the length of vector (maximum: %d) n=", MAX_N);
  scanf("%d",&n);
  printf("Enter the max value of vector elements k=");
  scanf("%d",&k);
  printf("Enter the sum of vector elements sum=");
  scanf("%d",&sum);
  int count = 0;
  gettimeofday(&t1, NULL);
  k++;

  thrust::host_vector<unsigned> test(n);
  thrust::fill(test.begin(), test.end(), 0);
  int finished = 0;
  do{
    if (thrust::reduce(test.begin(), test.end()) == sum) count++;
    finished = increment(test, k);
    }
    while (!finished);
  gettimeofday(&t2, NULL);
  printf("CPU count = %d, in %d seconds\n", count, t2.tv_sec - t1.tv_sec);
  unsigned long h_powers[MAX_N];
  h_powers[0] = 1;
  if (n < MAX_N)
    for (int i = 1; i<n; i++) h_powers[i] = h_powers[i-1]*k;
  cudaMemcpyToSymbol(powers, h_powers, MAX_N*sizeof(unsigned long));
  cudaCheckErrors("cudaMemcpyToSymbolfail");
  unsigned long long int *h_count, *d_count;
  h_count = (unsigned long long int *)malloc(sizeof(unsigned long long int));
  cudaMalloc((void **)&d_count, sizeof(unsigned long long int));
  cudaCheckErrors("cudaMalloc fail");
  *h_count = 0;
  cudaMemcpy(d_count, h_count, sizeof(unsigned long long int), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy H2D fail");
  find_vector_match<<<(GRIDSIZE + nTPB -1)/nTPB, nTPB>>>(d_count, k, n, sum);
  cudaMemcpy(h_count, d_count, sizeof(unsigned long long int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy D2H fail");
  gettimeofday(&t3, NULL);
  printf("GPU count = %d, in %d seconds\n", *h_count, t3.tv_sec - t2.tv_sec);

  return 0;
}

编译：

$ nvcc -O3 -arch=sm_20 -o t260 t260.cu

样本输出：

$ ./t260
Enter the length of vector (maximum: 12) n=2
Enter the max value of vector elements k=3
Enter the sum of vector elements sum=4
CPU count = 3, in 0 seconds
GPU count = 3, in 0 seconds
$ ./t260
Enter the length of vector (maximum: 12) n=5
Enter the max value of vector elements k=3
Enter the sum of vector elements sum=10
CPU count = 101, in 0 seconds
GPU count = 101, in 0 seconds
$ ./t260
Enter the length of vector (maximum: 12) n=9
Enter the max value of vector elements k=9
Enter the sum of vector elements sum=20
CPU count = 2714319, in 12 seconds
GPU count = 2714319, in 1 seconds
$ ./t260
Enter the length of vector (maximum: 12) n=10
Enter the max value of vector elements k=9
Enter the sum of vector elements sum=20
CPU count = 9091270, in 123 seconds
GPU count = 9091270, in 4 seconds

因此，对于较大的问题规模，朴素的蛮力 GPU 代码似乎比朴素的蛮力单线程 CPU 代码快约 30 倍。（...在我的特定机器设置上：CPU = Xeon X5560，GPU = Quadro5000，CentOS 5.5，CUDA 5.0）

于 2013-10-19T05:36:09.897 回答
0

问题是__syncthreads()。为了使__syncthreads()正常工作，块内的所有线程都应该能够访问它，否则一些线程将永远等待并且您的程序不会退出。在您的程序中，在某些部分执行__syncthreads()是有条件的。这就是为什么您的程序不能在一个块中使用多个线程的原因。

于 2013-10-18T19:20:22.860 回答
0

正如罗伯特在评论中所说，如果你想在 GPU 上生成所有 (k+1)^n 排列并测试它们，你可以考虑这样的一些 GPU 内核：

__device__ int count;  //global variable must be initialized to zero before kernel call
__global__ void perm_generator(int k, int n, int sum) {
   int tid = blockIdx.x*blockDim.x+threadIdx.x;
   int id = tid;
   int mysum = 0;
   for ( int i = n; i > 1; i-- ) { //all n-1 vector elements
     mysum += (id % (k+1));
     id /= (k+1);
   }
   mysum += id; //last element
   if ( mysum == sum ) atomicAdd( &count, 1 );
}

内核应该被精确地调用 (k+1)^n 个线程。如果你碰巧用更多线程调用你的内核（因为经验法则是块尺寸应该是 32 的倍数），你需要事先检查内核中tid的值。此外，不推荐使用cudaThreadSynchronize()。请改用cudaDeviceSynchronize()。

于 2013-10-19T00:45:33.560 回答