在CUDA中将char*转换为unsigned int*

Question

我有一个类型定义为typedef unsigned char uint40[5],然后我有一个uint40数组,比方说uint40* payloads

我试图将以下函数移植到CUDA内核中

void aSimpleFunction(int M, uint40* data)
{
    for (auto i = 0; i < M; i++)
    {
        unsigned int* dataPtr = (unsigned int*)data[i];
        *dataPtr = 2158677232;
        data[i][4] = 1;
    }
}

对我来说这很简单,但它没有用.但是,使用方括号来访问它确实有效的每个元素.

__global__ void aSimpleKernel(int M, uint40* data)
{
    int tid = threadIdx.x + 1;

    // DOESN'T WORK
    unsigned int* dataPtr = (unsigned int*)data[tid];
    *dataPtr = 16976944;
    // WORKS
    /*
    data[threadIdx.x][0] = tid * 1;
    data[threadIdx.x][1] = tid * 2;
    data[threadIdx.x][2] = tid * 3;
    data[threadIdx.x][3] = tid * 4;
    */
    data[threadIdx.x][4] = 2;
}

是否可以将char*转换为CUDA内核中的unsigned int*？

通过"不起作用"我的意思是,它具有随机数而不是我在打印uint40*数组的每个元素时所期望的.有时,GPU显然崩溃了,因为窗口中有一个弹出窗口告诉我gpu重启成功.

Answer 1

每当您遇到CUDA代码时遇到问题,最好使用正确的cuda错误检查并运行代码cuda-memcheck.即使你不理解错误输出,它对那些试图帮助你的人也很有用,所以我建议你在这里寻求帮助之前这样做.

我试图用你所展示的东西制作一个完整的代码是这样的:

#include <stdio.h>

typedef unsigned char uint40[5];


void aSimpleFunction(int M, uint40* data)
{
    for (int i = 0; i < M; i++)
    {
        unsigned int* dataPtr = (unsigned int*)data[i];
        *dataPtr = 0x02020202U;
        data[i][4] = 1;
    }
}

void uint40_print(uint40 &data){

  char *my_data = (char *)&data;
  for (int i = 0; i < 5; i++) printf("%d", my_data[i]);
  printf("\n");
}

__global__ void aSimpleKernel(int M, uint40* data)
{
    for (int i = 0; i < M; i++)
    {
        unsigned int* dataPtr = (unsigned int*)data[i];
        printf("%p\n", dataPtr);
        *dataPtr = 0x02020202U;
        data[i][4] = 1;
    }
}

int main(){

  uint40 *payloads = (uint40 *)malloc(10000);
  memset(payloads, 0, 10000);
  aSimpleFunction(5, payloads);
  uint40_print(payloads[0]);
  memset(payloads, 0, 10000);
  uint40 *d_payloads;
  cudaMalloc(&d_payloads, 10000);
  aSimpleKernel<<<1,1>>>(5, d_payloads);
  cudaMemcpy(payloads, d_payloads, 10000, cudaMemcpyDeviceToHost);
  for (int i = 0; i < 5; i++) uint40_print(payloads[i]);
  return 0;
}

当我编译并运行该代码时,我得到如下输出:

$ ./t1091
22221
00000
$

果然,GPU输出与CPU输出不匹配.如果我运行代码cuda-memcheck,我得到的输出的一部分看起来像这样:

$ cuda-memcheck ./t1091
========= CUDA-MEMCHECK
22221
========= Invalid __global__ write of size 4
=========     at 0x00000080 in /home/bob/misc/t1091.cu:28:aSimpleKernel(int, unsigned char[5]*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x402500005 is misaligned

这给出了实际问题的线索.实际上,您正在创建一个char数组,然后在其上叠加一个5字节宽的结构(uint40).这意味着连续uint40项将从相差5的字节地址开始.

当你把其中的一个地址,并把它转换到一个int或unsigned int指针,你可以用一个指针错位结束.CUDA要求POD数据类型的所有访问都在自然对齐的边界上发生.因此,必须在4字节边界(0,4,8,...)上访问32位数量(例如int,float等).uint40(0,5,10,...)的许多5字节边界也不落在4字节边界上,因此尝试以这种方式访问​​4字节数量是非法的.

一个可能的解决方案,对于这个特定的用法示例,并假设您传递给内核的指针是一个指针cudaMalloc(由于对齐),只是为了更改你的typedef:

typedef unsigned char uint40[8];

这会强制每个uint40项目落在一个8字节的边界上,这也是一个4字节的边界.这样做的副作用是分配每8个未使用的字节.

在您的情况下,您指出该uint40类型是数据的集合,而不是单个数字量,因此它实际上是一个数据"结构",恰好占用每个元素5个字节.这种"结构"的阵列实际上将是AoS(结构阵列)存储格式,并且对这种数据的性能的共同转换是将其转换为SoA(阵列结构)存储格式.因此,另一种可能的方法是创建两个数组:

typedef unsigned char uint40a[4];
typedef unsigned char uint40b[1];
uint40a *data1;
uint40b *data2;
cudaMalloc(&data1, size);
cudaMalloc(&data2, size);

并以这种方式访问​​您的数据.与5字节结构相比,这将保持存储密度,几乎可以肯定地提供对GPU中数据的更快访问.

如果从上面有任何疑问,你不能拿起一个任意char指针,把它转换成另一个(更大的)数据类型,并期望好事发生.您使用的指针必须与所引用的数据类型正确对齐.