让我们假设我们需要将以下字符串存储在 CUDA 数组中。
“你好呀”
“这是”
“谁是”
我们如何在 GPU 上声明一个数组来做到这一点。我尝试使用 C++ strings
,但它不起作用。
可能最好的方法是使用类似于常见压缩稀疏矩阵格式的结构。将打包的字符数据存储到一块线性内存中,然后使用一个单独的整数数组来存储起始索引,并可能使用第三个数组来存储字符串长度。后者的存储开销可能比为数据中的每个条目存储一个字符串终止字节并尝试在 GPU 代码中解析终止符更有效。
所以你可能有这样的事情:
struct gpuStringArray {
unsigned int * pos;
unsigned int * length; // could be a smaller type if strings are short
char4 * data; // 32 bit data type will improve memory throughput, could be 8 bit
}
注意我char4
为字符串数据使用了一个类型;矢量类型将提供更好的内存吞吐量,但这意味着字符串需要对齐/适当地填充到 4 字节边界。这可能是也可能不是问题,具体取决于您的应用程序中典型的真实字符串的外观。此外,可能应该选择(可选)长度参数的类型以反映最大允许字符串长度。如果您有很多非常短的字符串,则可能值得使用 8 位或 16 位无符号类型的长度来节省内存。
比较以这种方式存储的字符串的非常简单的代码可能如下strcmp
所示:
__device__ __host__
int cmp4(const char4 & c1, const char4 & c2)
{
int result;
result = c1.x - c2.x; if (result !=0) return result;
result = c1.y - c2.y; if (result !=0) return result;
result = c1.z - c2.z; if (result !=0) return result;
result = c1.w - c2.w; if (result !=0) return result;
return 0;
}
__device__ __host__
int strncmp4(const char4 * s1, const char4 * s2, const unsigned int nwords)
{
for(unsigned int i=0; i<nwords; i++) {
int result = cmp4(s1[i], s2[i]);
if (result != 0) return result;
}
return 0;
}
__global__
void tkernel(const struct gpuStringArray a, const gpuStringArray b, int * result)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
char4 * s1 = a.data + a.pos[idx];
char4 * s2 = b.data + b.pos[idx];
unsigned int slen = min(a.length[idx], b.length[idx]);
result[idx] = strncmp4(s1, s2, slen);
}
[免责声明:从未编译,从未测试,没有真实或暗示的保证,使用风险自负]
其中有一些极端情况和假设可能会根据代码中的真实字符串的确切外观来吸引您,但我会将这些作为练习留给读者解决。您应该能够适应并将其扩展到您正在尝试做的任何事情。
您必须使用 C 风格的字符串char *str
。在谷歌上搜索“CUDA string”会给你这个 CUDA“Hello World”示例作为第一次点击:http ://computer-graphics.se/hello-world-for-cuda.html
在那里你可以看到如何使用char*
- CUDA 中的字符串。请注意,标准 C 函数在 CUDA 中可用strcpy
或strcmp
不可用!
如果你想要一个字符串数组,你只需要使用char**
(如在 C/C++ 中)。至于strcmp
和类似的功能,这在很大程度上取决于你想做什么。CUDA 不太适合字符串操作,如果您能提供更多关于您想要做什么的细节,也许会有所帮助。