1
1
运行时
CUDA C++ block设置导致运行结果不一致
2023-09-19 11:38:25
121

运行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
全部评论(2)
只看楼主
T
tianlh
2楼

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
2023/9/25
L
lizhigong
1楼
gbz = Nz / btz;这一句计算结果是0,会造成kernel发起失败,你可以在launch后面加上cudaGetLastError();获取一下发起的结果。
2023/9/20
没有更多内容了
--
楼主
Level.1
0 帖子
楼主热帖
暂无数据