通过更改线程数更改 CUDA 代码输出的说明
Posted
技术标签:
【中文标题】通过更改线程数更改 CUDA 代码输出的说明【英文标题】:Explanation for change in output of CUDA code by changing number of threads 【发布时间】:2013-07-06 10:40:34 【问题描述】:我想确定有多少 x^2+1 形式的数是素数,因为 1
我布置了一个网格,并在我的区间内滑动,将结果记录在每个块的共享内存中,对每个块执行 gpu 缩减,最后执行 cpu 缩减以获得最终结果。
我的问题是,当我更改块数和每个块中的线程数时,输出结果会发生变化。我无法解释的另一件事是,对于 8 个块和每个块 2048 个线程的配置,代码在 100 毫秒内运行,但是当我将线程数减少到 1024 并将块数加倍时,代码将导致超时在从设备到主机的 memcpy 中!!我该如何解释这种行为以及正确性出现问题的地方?
我使用的是 GTX 480 nvidia gpu。
我的代码是:
#include <stdio.h>
static void HandleError( cudaError_t err, const char *file, int line )
if (err != cudaSuccess)
printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
exit( EXIT_FAILURE );
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
#define N 10000000
#define BLOCKS 8
#define THREADS 2048
__device__ int isprime(int x)
long long n = (long long)x*x + 1;
for( int p=3; p<=x+1; p+=2 )
if ( n % p == 0 ) return 0;
return 1;
__global__ void solve(int n, int* result)
__shared__ int ipc[THREADS];
int tid = threadIdx.x;
int x = blockIdx.x*blockDim.x + threadIdx.x + 2;
// sliding grid window over interval of to-be-computed data
int acc = 0;
while( x <= n )
if ( isprime(x) ) acc++;
x += blockDim.x*gridDim.x;
ipc[tid] = acc;
__syncthreads();
// reduction over each block in parallel
for( int s=blockDim.x/2; s>0; s>>=1 )
if ( tid < s )
ipc[tid] += ipc[tid+s];
__syncthreads();
if ( tid == 0 ) result[blockIdx.x] = ipc[0];
int main()
int *dev;
int res[BLOCKS];
int ans = 0;
HANDLE_ERROR( cudaMalloc((void**)&dev, BLOCKS * sizeof(int)) );
solve<<<BLOCKS, THREADS>>>(N, dev);
HANDLE_ERROR( cudaMemcpy(res, dev, BLOCKS*sizeof(int), cudaMemcpyDeviceToHost) );
// final reduction over results for each block
for( int j=0; j<BLOCKS; j++ )
ans += res[j];
printf("ans = %d\n", ans);
HANDLE_ERROR( cudaFree( dev ) );
return 0;
【问题讨论】:
【参考方案1】:您不能在任何当前 GPU 上每个块运行 2048 个线程:
#define THREADS 2048
...
solve<<<BLOCKS, THREADS>>>(N, dev);
^
|
2048 is illegal here
您没有在内核调用上正确地执行cuda error checking,因此您的代码不会告诉您正在发生此错误。
因此,在每个块有 2048 个线程的情况下,您的内核甚至没有执行(而且您的结果应该是虚假的。)
如果你将线程减半,超时可能是由于你的内核执行时间过长,windows TDR mechanism 启动了。
我尝试使用 BLOCKS
= 16 和 THREADS
= 1024 运行您的代码
在 N = 100000 时,我的 M2050 GPU 上的总执行时间约为 1.5 秒。 N = 1000000 时,执行时间约为 75 秒。 N = 10000000 这就是你所拥有的,执行时间非常长。
【讨论】:
以上是关于通过更改线程数更改 CUDA 代码输出的说明的主要内容,如果未能解决你的问题,请参考以下文章