c++ - 如何理解将半精度指针转换为无符号long long指针以及相关的内存对齐方式？

Question

我不熟悉有关内存对齐和指针转换的问题。我正在从Nvidia官方示例代码中学习以下内容。

  half *A = NULL;
  half *B = NULL;
  float *C = NULL;
  float *D = NULL;

  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&A),
                             sizeof(half) * M_GLOBAL * K_GLOBAL));
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&B),
                             sizeof(half) * N_GLOBAL * K_GLOBAL));
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&C),
                             sizeof(float) * M_GLOBAL * N_GLOBAL));
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&D),
                             sizeof(float) * M_GLOBAL * N_GLOBAL));

  assert(((unsigned long long)A) % 128 == 0);
  assert(((unsigned long long)B) % 128 == 0);
  assert(((unsigned long long)C) % 128 == 0);
  assert(((unsigned long long)D) % 128 == 0);

我对此有三个问题。

我知道指针类型转换和数值类型转换之间的区别。关于指针类型转换，我们只是改变了理解该指针的方式。但是我应该如何理解从half*到unsigned long long*的转换？

我认为此转换期间的内存地址应保持不变。为什么我们需要先转换为unsigned long long*，然后检查内存是否与128对齐？

Answer 1

But how should I understand the conversion from half* to unsigned long long*?

您显示的代码中没有 unsigned long long *的转换。有一个到 unsigned long long的转换。

转换的目的是将 A， B， C或 D之一中存储的地址转换为整数，以便可以检查其位。 C标准未定义将指针转换为整数类型的结果(除某些基本属性外)，但该转换“旨在与执行环境的寻址结构一致”(C 2018脚注69)。在Nvidia使用的编译器中，转换大概会产生处理器架构通常使用的地址。然后使用 % 128 == 0测试地址是否对齐128字节的倍数。

Why do we need to convert to unsigned long long* first and then check if the memory is aligned with 128?

%运算符将不接受指针操作数，因此必须将操作数转换为整数类型 unsigned long long，而不是 unsigned long long *。