为什么简单的 CUDA 程序的结果每次都不同?
Why the result of simple CUDA program differs every time?
想做一个简单的tiling卷积码。来自Coursera :Heterogeneous Parallel Programming 的讲座。讲座提供了一个简单的tiling方法的卷积代码,但是代码并不完整。因此,我填写了代码中的空白,下面是结果。
这段代码的目标是计算卷积。
输入尺寸:(24 x 24),
内核大小:(9 x 9),
输出尺寸:(16 x 16)。
另外,主要是,我放了计算时间检查代码来与CPU版本进行比较。
问题是,每当我 运行 这段代码时,结果都是不同的。找了好几天的问题,每次都试不出来。我也在网上找到了类似的代码 blog 但它和我有同样的问题。不知道为什么每次的结果都不一样。有人说这是由于比赛条件造成的,但我没有找到任何相关信息。
这是卷积的示例结果(16 x 16 大小)。
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
0 0 0 0 0 0 0 0 0 0 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
我的设备是 CUDA Driver = CUDART,CUDA Driver Version = 7.5,CUDA Runtime Version = 7.5,NumDevs = 1,Device0 = GeForce GT 630。
我使用 Ubuntu 14.04.
提前致谢!
#include<stdio.h>
#include<cuda.h>
#include<time.h>
#define O_TILE_WIDTH 10
#define MASK_WIDTH 9
#define I_TILE_WIDTH (O_TILE_WIDTH+MASK_WIDTH-1)
__global__ void Convolution2DBasicKernel(float *out, float *in, int in_height, int in_width, const float *__restrict__ mask, int output_dim)
{
int tx=threadIdx.x;
int ty=threadIdx.y;
int row_o=blockIdx.y*O_TILE_WIDTH+ty;
int col_o=blockIdx.x*O_TILE_WIDTH+tx;
int row_i=row_o;
int col_i=col_o;
__syncthreads();
__shared__ float Ns[I_TILE_WIDTH][I_TILE_WIDTH];
///////////////////////////////////////////////////////////
//////////////////// reading input data ///////////////////
if( (row_i>=0)&&(row_i<in_height)&&(col_i>=0)&&(col_i<in_width) )
{
Ns[ty][tx]=in[row_i*in_width + col_i];
}
else
{
Ns[ty][tx]=0.0f;
}
__syncthreads();
///////////////////////////////////////////////////////////
//////////////////// calculating convol ///////////////////
float output=0.0f;
if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) )
{
for(int i=0; i<MASK_WIDTH; i++)
{
for(int j=0; j<MASK_WIDTH; j++)
{
output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];
}
}
}
__syncthreads();
if( (row_o<output_dim)&&(col_o<output_dim) )
{
out[row_o*output_dim+col_o]=output;//in_width
}
__syncthreads();
}
int main() {
int input_dim=24;
int kernel_dim=9;
int output_dim=16;
float *input = new float[input_dim*input_dim];
float *kernel = new float[kernel_dim*kernel_dim];
float *output = new float[output_dim*output_dim];
float *d_input;
float *d_kernel;
float *d_output;
cudaMalloc(&d_input, sizeof(float)*input_dim*input_dim);
cudaMalloc(&d_kernel, sizeof(float)*kernel_dim*kernel_dim);
cudaMalloc(&d_output, sizeof(float)*output_dim*output_dim);
for(int i=0; i<input_dim*input_dim; i++)
{
input[i]=1.0;
}
for(int i=0; i<kernel_dim*kernel_dim; i++)
{
kernel[i]=1.0;
}
cudaMemcpy(d_input, input, sizeof(float)*input_dim*input_dim, cudaMemcpyHostToDevice);
cudaMemcpy(d_kernel, kernel, sizeof(float)*kernel_dim*kernel_dim, cudaMemcpyHostToDevice);
dim3 dimBlock (I_TILE_WIDTH, I_TILE_WIDTH, 1);
dim3 dimGrid ((output_dim-1)/O_TILE_WIDTH+1, (output_dim-1)/O_TILE_WIDTH+1, 1);
clock_t begin, end;
double time_spent;
begin = clock();
for(int iteration=0; iteration<1; iteration++)//100000
{
Convolution2DBasicKernel<<<dimGrid, dimBlock>>>(d_output, d_input, input_dim, input_dim, d_kernel, output_dim);
}
end = clock();
time_spent = (double)(end - begin) / CLOCKS_PER_SEC;
printf("time: %f\n", time_spent);
cudaMemcpy(output, d_output, sizeof(float)*output_dim*output_dim, cudaMemcpyDeviceToHost);
for(int y=0; y<output_dim; y++)
{
for(int x=0; x<output_dim; x++)
printf("%d\t", int(output[y*16+x]));
printf("\n");
}
}
您正在越界访问共享内存。
假设您确信您的程序或多或少是正确的,您将需要确保您没有越界:
if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) ) {
for(int i=0; i<MASK_WIDTH; i++) {
if(ty +i < O_TILE_WIDTH) { // Changed here
for(int j=0; j<MASK_WIDTH; j++) {
if(tx +j < O_TILE_WIDTH) { // Changed here
output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];
}
}
}
}
}
这是一个竞争条件。
这是一个例子。
您正在为每个块和 2x2 块启动 18x18 线程。
示例:
ThreadA
threadIdx.x = 10
threadIdx.y = 0
blockIdx.x = 0
blockIdx.y = 0
ThreadB
threadIdx.x = 0
threadIdx.y = 0
blockIdx.x = 1
blockIdx.y = 0
计算时在内核内部:
int tx=threadIdx.x
int ty=threadIdx.y
int row_o=blockIdx.y*O_TILE_WIDTH+ty
int col_o=blockIdx.x*O_TILE_WIDTH+tx
using O_TILE_WIDTH = 10
ThreadA
row_o = 0*10+10 = 10
col_o = 0
ThreadB
row_o = 1*10+0 = 10
col_o = 0
这意味着两个线程将在内存中的相同位置输出结果,但计算方式不同。
想做一个简单的tiling卷积码。来自Coursera :Heterogeneous Parallel Programming 的讲座。讲座提供了一个简单的tiling方法的卷积代码,但是代码并不完整。因此,我填写了代码中的空白,下面是结果。
这段代码的目标是计算卷积。
输入尺寸:(24 x 24),
内核大小:(9 x 9),
输出尺寸:(16 x 16)。
另外,主要是,我放了计算时间检查代码来与CPU版本进行比较。
问题是,每当我 运行 这段代码时,结果都是不同的。找了好几天的问题,每次都试不出来。我也在网上找到了类似的代码 blog 但它和我有同样的问题。不知道为什么每次的结果都不一样。有人说这是由于比赛条件造成的,但我没有找到任何相关信息。
这是卷积的示例结果(16 x 16 大小)。
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
0 0 0 0 0 0 0 0 0 0 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 81 81 81 81 81 81
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
81 81 81 81 81 81 81 81 81 81 0 0 0 0 0 0
我的设备是 CUDA Driver = CUDART,CUDA Driver Version = 7.5,CUDA Runtime Version = 7.5,NumDevs = 1,Device0 = GeForce GT 630。 我使用 Ubuntu 14.04.
提前致谢!
#include<stdio.h>
#include<cuda.h>
#include<time.h>
#define O_TILE_WIDTH 10
#define MASK_WIDTH 9
#define I_TILE_WIDTH (O_TILE_WIDTH+MASK_WIDTH-1)
__global__ void Convolution2DBasicKernel(float *out, float *in, int in_height, int in_width, const float *__restrict__ mask, int output_dim)
{
int tx=threadIdx.x;
int ty=threadIdx.y;
int row_o=blockIdx.y*O_TILE_WIDTH+ty;
int col_o=blockIdx.x*O_TILE_WIDTH+tx;
int row_i=row_o;
int col_i=col_o;
__syncthreads();
__shared__ float Ns[I_TILE_WIDTH][I_TILE_WIDTH];
///////////////////////////////////////////////////////////
//////////////////// reading input data ///////////////////
if( (row_i>=0)&&(row_i<in_height)&&(col_i>=0)&&(col_i<in_width) )
{
Ns[ty][tx]=in[row_i*in_width + col_i];
}
else
{
Ns[ty][tx]=0.0f;
}
__syncthreads();
///////////////////////////////////////////////////////////
//////////////////// calculating convol ///////////////////
float output=0.0f;
if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) )
{
for(int i=0; i<MASK_WIDTH; i++)
{
for(int j=0; j<MASK_WIDTH; j++)
{
output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];
}
}
}
__syncthreads();
if( (row_o<output_dim)&&(col_o<output_dim) )
{
out[row_o*output_dim+col_o]=output;//in_width
}
__syncthreads();
}
int main() {
int input_dim=24;
int kernel_dim=9;
int output_dim=16;
float *input = new float[input_dim*input_dim];
float *kernel = new float[kernel_dim*kernel_dim];
float *output = new float[output_dim*output_dim];
float *d_input;
float *d_kernel;
float *d_output;
cudaMalloc(&d_input, sizeof(float)*input_dim*input_dim);
cudaMalloc(&d_kernel, sizeof(float)*kernel_dim*kernel_dim);
cudaMalloc(&d_output, sizeof(float)*output_dim*output_dim);
for(int i=0; i<input_dim*input_dim; i++)
{
input[i]=1.0;
}
for(int i=0; i<kernel_dim*kernel_dim; i++)
{
kernel[i]=1.0;
}
cudaMemcpy(d_input, input, sizeof(float)*input_dim*input_dim, cudaMemcpyHostToDevice);
cudaMemcpy(d_kernel, kernel, sizeof(float)*kernel_dim*kernel_dim, cudaMemcpyHostToDevice);
dim3 dimBlock (I_TILE_WIDTH, I_TILE_WIDTH, 1);
dim3 dimGrid ((output_dim-1)/O_TILE_WIDTH+1, (output_dim-1)/O_TILE_WIDTH+1, 1);
clock_t begin, end;
double time_spent;
begin = clock();
for(int iteration=0; iteration<1; iteration++)//100000
{
Convolution2DBasicKernel<<<dimGrid, dimBlock>>>(d_output, d_input, input_dim, input_dim, d_kernel, output_dim);
}
end = clock();
time_spent = (double)(end - begin) / CLOCKS_PER_SEC;
printf("time: %f\n", time_spent);
cudaMemcpy(output, d_output, sizeof(float)*output_dim*output_dim, cudaMemcpyDeviceToHost);
for(int y=0; y<output_dim; y++)
{
for(int x=0; x<output_dim; x++)
printf("%d\t", int(output[y*16+x]));
printf("\n");
}
}
您正在越界访问共享内存。
假设您确信您的程序或多或少是正确的,您将需要确保您没有越界:
if( (tx<O_TILE_WIDTH)&&(ty<O_TILE_WIDTH) ) {
for(int i=0; i<MASK_WIDTH; i++) {
if(ty +i < O_TILE_WIDTH) { // Changed here
for(int j=0; j<MASK_WIDTH; j++) {
if(tx +j < O_TILE_WIDTH) { // Changed here
output += Ns[ty+i][tx+j]*mask[i*MASK_WIDTH+j];
}
}
}
}
}
这是一个竞争条件。
这是一个例子。
您正在为每个块和 2x2 块启动 18x18 线程。
示例:
ThreadA
threadIdx.x = 10
threadIdx.y = 0
blockIdx.x = 0
blockIdx.y = 0
ThreadB
threadIdx.x = 0
threadIdx.y = 0
blockIdx.x = 1
blockIdx.y = 0
计算时在内核内部:
int tx=threadIdx.x
int ty=threadIdx.y
int row_o=blockIdx.y*O_TILE_WIDTH+ty
int col_o=blockIdx.x*O_TILE_WIDTH+tx
using O_TILE_WIDTH = 10
ThreadA
row_o = 0*10+10 = 10
col_o = 0
ThreadB
row_o = 1*10+0 = 10
col_o = 0
这意味着两个线程将在内存中的相同位置输出结果,但计算方式不同。