4位整型数组的内存布局

Question

int4和int2数组的内存布局是怎样的？认为，

int4 M[2];
M[0]=0xA;
M[1]=0x5;

我应该在包含 M(0) 和 M(1) 的单个字节中看到什么？是0xA5还是0x5A？我知道 int4 不是 C/C++ 类型，但某些编译器必须处理它，因为 Nvidia 和 AMD GPU 支持它。

Answer 1

CUDA 中的数据类型int4由 4 个 32 位整数组成。

然而，NVidia 的张量核心 wmma 操作可以处理子字节元素，例如打包在 32 位整数中的 4 位整数和 1 位整数，但这些子字节元素无法通过[]. Hopper (sm_90) 中已弃用并删除了 4 位操作。编程指南中对此进行了解释。 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#sub-byte-operations

子字节 WMMA 运算提供了一种访问 Tensor Core 的低精度功能的方法。它们被视为预览功能，即它们的数据结构和 API 可能会发生变化，并且可能与未来版本不兼容。此功能可通过 nvcuda::wmma::experimental 命名空间使用：

namespace experimental {
    namespace precision {
        struct u4; // 4-bit unsigned
        struct s4; // 4-bit signed
        struct b1; // 1-bit
   }
...
}

以下代码打印 wmma 片段的内容，其原始数据的第一个字节设置为0b10100101(0xA5)

#include <iostream>
#include <vector>

#include <mma.h>
using namespace nvcuda;

__global__ void kernel(unsigned int *a, unsigned int *b, unsigned int *c) {
    wmma::fragment<wmma::matrix_a, 8, 8, 32, nvcuda::wmma::experimental::precision::u4, wmma::row_major> a_frag;

    wmma::load_matrix_sync(a_frag, a, 32);


    for(int t = 0; t < 1; t++){
        if(threadIdx.x == t){
            printf("thread %d, fragment.num_elements %d\n", threadIdx.x, a_frag.num_elements);
            for(int t=0; t < a_frag.num_elements; t++){
                int v = a_frag.x[t];
                printf("%d ", v);
            }
            printf("\n");
        }
    }
}

int main(){
    unsigned int* d_A; cudaMalloc(&d_A, sizeof(unsigned int) * 1024);
    unsigned int* d_B; cudaMalloc(&d_B, sizeof(unsigned int) * 1024);
    unsigned int* d_C; cudaMalloc(&d_C, sizeof(unsigned int) * 1024);

    std::vector<unsigned char> bytes(1024);
    bytes[0] = 0b10100101;
    cudaMemcpy(d_A, bytes.data(), 1024, cudaMemcpyHostToDevice);

    kernel<<<1, 32>>>(d_A, d_B, d_C);
    cudaDeviceSynchronize();
}

165该代码10100101以二进制形式打印。