即使未在主机代码中初始化步幅，cuda 内核也会自动为矩阵设置步幅吗？

Question

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

这是矩阵的结构


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 编程指南中第 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 并在此设备代码中 "which depends only on device memory" 您会发现这些行：

// 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 中的结构一样从矩阵中扣除 A.stride ..
但是在 2 个不同矩阵的主机代码和设备代码中使用名称 A 导致我感到困惑。

问题已解决。

Answer 1

在 GetSubMatrix 的设备代码中使用名称 A 来引用主机代码矩阵并在设备代码中引用 d_A 矩阵导致混淆，因为 struct Matrix 的数据成员步长未在主机代码矩阵中初始化，但它将在设备副本d_A矩阵中初始化，
并且此 d_A 将通过名为 A 的参数传递给 GetSubMatrix，该参数定义了步幅。
所以我们有 2 个矩阵，名称为 A，一个在主机中未定义，另一个在设备中已定义，所以我有这种误解。

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

即使未在主机代码中初始化步幅，cuda 内核也会自动为矩阵设置步幅吗？

do cuda kernel sets stride for matrix automatically even if stride not initialized in host code?

c++

matrix

cuda

stride