运行CUDA C++的时候,不同线程块配置导致运行结果不一致,具体测试代码如下:
#include <iostream>
#include <string>
#include <vector>
#include <cuda.h>
#include "cuda_runtime.h"
#define INDEX(i, j, k) \
(__mul24((i), __mul24((Ny), (Nz))) + __mul24((j), (Nz)) + (k))
// add two vector of same size: c[i] = a[i] + b[i]
__global__ void add(float *a, float *b, float *c, int Nx, int Ny, int Nz)
{
int nx = blockIdx.x * blockDim.x + threadIdx.x;
int ny = blockIdx.y * blockDim.y + threadIdx.y;
int nz = blockIdx.z * blockDim.z + threadIdx.z;
if (nx < Nx && ny < Ny && nz < Nz)
{
c[INDEX(nx, ny, nz)] = a[INDEX(nx, ny, nz)] + b[INDEX(nx, ny, nz)];
}
}
float test(int _Nx, int _Ny, int _Nz, int _btx, int _bty, int _btz)
{
int Nx, Ny, Nz;
Nx = _Nx;
Ny = _Ny;
Nz = _Nz;
// block threads pattern: btx = 1, bty = 1, btz = 256
// grid blocks pattern: gbx = Nx/btx, gby = Ny/bty, gbz = Nz/btz
int btx, bty, btz;
btx = _btx;
bty = _bty;
btz = _btz;
int gbx, gby, gbz;
gbx = Nx / btx;
gby = Ny / bty;
gbz = Nz / btz;
int size = Nx * Ny * Nz;
float *a, *b, *c;
cudaMallocManaged(&a, size * sizeof(float));
cudaMallocManaged(&b, size * sizeof(float));
cudaMallocManaged(&c, size * sizeof(float));
for (int i = 0; i < size; i++)
{
a[i] = 1.0f;
b[i] = 2.0f;
}
dim3 block(btx, bty, btz);
dim3 grid(gbx, gby, gbz);
add<<<grid, block>>>(a, b, c, Nx, Ny, Nz);
cudaDeviceSynchronize();
float maxError = 0.0f;
for (int i = 0; i < size; i++)
{
maxError = fmax(maxError, fabs(c[i] - 3.0f));
}
return maxError;
}
int main()
{
{
// Case 1: Nx = 128, Ny = 128, Nz = 128, btx = 1, bty = 1, btz = 128
std::string casename = "Case 1";
int Nx, Ny, Nz;
Nx = 128;
Ny = 128;
Nz = 128;
int btx, bty, btz;
btx = 1;
bty = 1;
btz = 128;
float maxError = test(Nx, Ny, Nz, btx, bty, btz);
std::cout << casename << " max error: " << maxError << std::endl;
}
{
// Case 2: Nx = 128, Ny = 128, Nz = 128, btx = 1, bty = 1, btz = 256
std::string casename = "Case 2";
int Nx, Ny, Nz;
Nx = 128;
Ny = 128;
Nz = 128;
int btx, bty, btz;
btx = 1;
bty = 1;
btz = 256;
float maxError = test(Nx, Ny, Nz, btx, bty, btz);
std::cout << casename << " max error: " << maxError << std::endl;
}
return 0;
}
加载的环境为:
module purge
module load compiler/dtk/23.04
source /public/software/compiler/rocm/dtk-23.04/env.sh
source /public/software/compiler/rocm/dtk-23.04/cuda/env.sh
module load compiler/cmake/3.23.1
module load compiler/gcc/9.3.0
运行结果为:
Case 1 max error: 0
Case 2 max error: 3

暂无数据
gbx = (Nx - 1) / btx + 1;
gby = (Ny - 1) / bty + 1;
gbz = (Nz - 1) / btz + 1;
请按照教程示例规范定义
https://developer.hpccube.com/gitbook//dcu_developer/DeveloperGuide/dcu_programming_2023/chapter3/3.1.4.html