即使跨度未在主机代码中初始化，cuda内核是不是会自动为矩阵设置跨度？ [关闭]

Posted 2023-03-23

【中文标题】即使跨度未在主机代码中初始化，cuda内核是不是会自动为矩阵设置跨度？ [关闭]

【英文标题】：do cuda kernel sets stride for matrix automatically even if stride not initialized in host code? [closed]即使跨度未在主机代码中初始化，cuda内核是否会自动为矩阵设置跨度？ [关闭]

【发布时间】：2019-11-23 10:43:30

【问题描述】：

我正在研究 cuda c，我正在使用的源代码使用 cuda 示例程序，特别是在运行时进行矩阵乘法。 我正在逐行跟踪代码并尝试预测下一步以确保我理解代码。 在此期间，我发现了具有数据成员 stride 的 Matrix 的结构声明。 整个代码没有单行初始化这个跨步数据成员。 我用nsight调试设备代码，用普通vs调试器调试主机代码>>>>>有惊喜： 直到程序成功结束，主机代码才真正初始化这个数据成员。 但 nsight 甚至在第一个内核行之前就显示了步幅已初始化。 当我查看调用内核的 vs 调试器的 autos 窗口时，我注意到内核的函数名称行显示 __cuda_0 矩阵，其结构与程序 Matrix 结构相同，但具有初始化步幅？？？？？？ 所以我不知道何时以及谁在设备代码上初始化了这个步幅变量？？？ 非常感谢

这是矩阵的结构

typedef struct 
   int width;
    int height;
    float* elements;    
    int stride;
  Matrix;

这是初始化矩阵的主要代码

int main(int argc, char* argv[])

    Matrix A, B, C;
    int a1, a2, b1, b2;

    a1 = atoi(argv[1]); /* Height of A */
    a2 = atoi(argv[2]); /* Width of A */
    b1 = a2; /* Height of B */
    b2 = atoi(argv[3]); /* Width of B */

    A.height = a1;
    A.width = a2;
    A.elements = (float*)malloc(A.width * A.height * sizeof(float));

    B.height = b1;
    B.width = b2;
    B.elements = (float*)malloc(B.width * B.height * sizeof(float));

    C.height = A.height;
    C.width = B.width;
    C.elements = (float*)malloc(C.width * C.height * sizeof(float));

    for(int i = 0; i < A.height; i++)
        for(int j = 0; j < A.width; j++)
            A.elements[i*A.width + j] = (rand() % 3);//arc4random

    for(int i = 0; i < B.height; i++)
        for(int j = 0; j < B.width; j++)
            B.elements[i*B.width + j] = (rand() % 2);//arc4random

    MatMul(A, B, C);

整个代码都存在于：CUDA C Programming Guide 第3-2-3章

好的，到目前为止我得到了-4，可能是问题的目的不明确： 在 MatMul 主机函数中，有几行声明和初始化所用矩阵的设备副本，它使用 A.width 来初始化 d_A.stride ....

 void MatMul(const Matrix A, const Matrix B, Matrix C) 

 // Load A and B to device memory
 Matrix d_A;
 d_A.width = d_A.stride = A.width;
 d_A.height = A.height;
 size_t size = A.width * A.height * sizeof(float);
 cudaMalloc(&d_A.elements, size);
 cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);

但是当你到达时：

 // Invoke kernel
 dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
 dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);  

它调用 MatMulKernel 并在此设备代码“仅取决于设备内存”中找到这些行：

// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)

它以矩阵 A 作为参数......在这里我看到了我的困惑原因！！！！ MatMulKernel 使用名称 A 来引用传递给它的 d_A 矩阵... 所以稍后在这些行：

    // Get sub-matrix Asub of A
    Matrix Asub = GetSubMatrix(A, blockRow, m);  

它调用另一个名为 GetSubMatrix 的设备函数，将 A 传递给它，它实际上是 d_A，然后在 GetSubMatrix 代码中它使用 A.stride，它实际上是 d_A.stride

__device__ Matrix GetSubMatrix(Matrix A, int row, int col) 

    Matrix Asub;
    Asub.width    = BLOCK_SIZE;
    Asub.height   = BLOCK_SIZE;
    ***Asub.stride   = A.stride;***
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row
                                     + BLOCK_SIZE * col];
    return Asub;
   

所以宿主代码结构确实没有初始化 A.stride 并且没有隐藏机制可以从 cuda 中的 structre 之类的矩阵中扣除 A.stride .. 但是在 2 个不同矩阵的主机代码和设备代码中使用名称 A 导致了我的困惑。

问题解决了。

【问题讨论】：

要查看的最重要的代码似乎是MatMul() 的源代码。果然，（任何人都可以检查online），有初始化：d_A.width = d_A.stride = A.width; d_A.height = A.height;。

是的，我看到了这段代码，这让我非常惊讶.... A.width 这里是 rhs 表达式，代码中没有其他地方可以初始化它...所以怎么可能它初始化 d_A.width 或 d_A.stride .........如果你尝试调试它，你会发现在主机代码中它保持未初始化直到程序结束......那么设备代码是如何获得了 A.stride 的价值？？？？？？

如果您仔细观察，您会注意到，虽然d_A.stride 确实出现在等号的右侧，但它也出现在另一个等号的左侧，即makes it a valid assignment。跨度>

【参考方案1】：

使用名称 A 来指代主机代码矩阵和在 GetSubMatrix 的设备代码中指代 d_A 矩阵会导致混淆，因为 struct Matrix 的数据成员跨度没有在主机代码矩阵中初始化，而是在设备中初始化复制 d_A 矩阵， 并且此 d_A 将通过名为 A 的参数传递给 GetSubMatrix，该参数已定义跨度。 所以我们有 2 个名称为 A 的矩阵，一个在主机中未定义，另一个在已定义的设备中，所以我有这个误解。

如果他们将 GetSubMatrix 中的参数名称从 A 更改为其他任何东西，则不会对 stride 数据成员造成混淆。