1

我在全局内存中有一个 int 数组。为了减少从全局内存中读取的频率，我一直在尝试使用 64 位数据类型进行读取，然后根据需要使用高 32 位或低 32 位。例如，这会从数组中获取第 3 和第 4 个整数：

__device__ void func1(int* arr)
{
    unsigned long long int val = *((unsigned long long int *) &arr[3]);

    // Now operate on the individual ints
}

使用这种方法来检索整数给了我未定义的行为，即使它看起来应该可行。当它起作用时，以这种方式读取值比单个整数读取要快得多。有没有人遇到过这个问题？

4

1 回答 1

1

数量喜欢按大小排列。我不确定 cuda 如何处理你正在做的事情以及它可能是特定于环境的，但是你使用：

*((unsigned long long int *) &arr[3])

假设arr是 8 字节对齐，则采用 8 字节数量，即只有 4 字节对齐。这当然是因为：

arr = 8n           // n is an integer
sizeof(int) = 4

&arr[3] = 8n + 3*4 // simplifies to 8(n+1) + 4

我知道如果您尝试在使用 32 位和 16 位整数的处理器上执行相同的操作（尽管我从未尝试过使用 64 位和 32 位整数），您会遇到问题。

您将需要自制某种访问器交易，以确定您尝试访问的数据在哪里。考虑以下情况，类似于您的情况：

int get32BitValueFrom(unsigned long long int longArray[], int index)
{
    // get the 64 bit int containing the 32 bit int we want
    unsigned long long int value = longarray[index >> 1];

    // if we wanted an odd index, return the high order 32 bits
    // otherwise return the low order 32 bits
    return (int) ((index & 1) ? (value >> 32) : (value));
}

编辑：我知道您正在使用 cuda，并且我知道要避免分支，但我确信有一种方法可以使用某种按位或数学技巧来编写等效代码来完成相同的事情。

于 2012-07-25T22:12:28.843 回答