【发布时间】：2013-10-19 13:36:09
【问题描述】：
我想在 CUDA C 中并行化一个函数，它将计算所有向量元素和不大于 k 的元素之和相等的向量。例如，如果向量元素的数量 n 为 5，sum=10 且 k=3，则满足此条件的向量数量为 101。我已经在 CUDA C 中制作了此函数，但问题是当块和线程大于 1。我知道问题出在 for 循环中，我应该更改它，但我不知道从哪里开始。当我使用等于 1 的块和线程调用函数时，该函数以经典方式工作，一切都很好，但在这种情况下，函数没有并行化。
程序源代码为：
//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;
}
【问题讨论】：
-
我假设向量元素的范围可以从0到
k，而k应该是一个正数？由于您使用 1 个线程的 1 个线程块调用内核，因此您根本没有真正并行化任何东西。您只是在运行串行代码。你有没有尝试过并行化？你甚至有策略吗？例如，如果你有多个线程，你会让每个线程做什么？ （也许是一个单独的模式测试？）每个块可能负责什么？ （也许是要测试的整体空间的一部分？） -
是的，k 是正数，我知道当块数和线程数为 1 时，代码是串行运行的。如何使用线程进行模式测试？
-
要测试的空间是所有长度为
n的向量，其中每个向量元素都有k+1个可能的值。因此，使用简单的蛮力方法，有 (k+1)^n 个可能的向量要测试。一种方法是让每个线程完成所有的工作（向量生成、总和计算、总和测试）以测试单个向量，然后让整个网格循环通过 (k+) 的整个空间1)^n 个向量。对于正确测试的向量，您可以在每个线程结束时使用原子操作来更新count值。这将适用于适合unsigned long的 (k+1)^n。 -
所以您建议生成所有可能的向量组合，并在生成后对一种组合使用一个线程并检查它是否是向量满足条件？
-
这是一种可能的方法。不是特别聪明，但它的好处是它应该很容易尝试，因为它类似于为简单的 CPU 串行实现编写代码的方式。而且我不会预先生成向量 - 你会用完空间来解决大问题。我会让每个线程使用基于线程索引 (
int idx = threadIdx.x + blockDim.x*blockIdx.x;) 的算法即时生成它的唯一向量，这样您一次只需要足够的存储空间来存储一个网格的向量。