PyCUDA 2D 数组实现（或使用字符串）

Question

我正在尝试在 CUDA 中使用字符串数组（words）。

我尝试通过创建单个字符串来将其展平，但随后要对其进行索引，每次内核运行时我都必须对其进行一些处理。如果有 9000 个字长为 6 个字符，那么在最坏的情况下，我必须检查每个内核调用的 53994 个字符。所以我正在寻找不同的方法来做到这一点。

更新：忘了说，字符串的长度不同，所以我必须找到每个字符串的结尾。

接下来我尝试将每个单词复制到不同的内存位置，然后收集地址，并将其作为数组传递给 GPU，代码如下：

# np = numpy

wordList = ['asd','bsd','csd']

d_words = []

for word in wordList:
    d_words.append(gpuarray.to_gpu(np.array(word, dtype=str)))

d_wordList = gpuarray.to_gpu(np.array([word.ptr for word in d_words], dtype=np.int32))

ker_test(d_wordList, block=(1,1,1), grid=(1,1,1))

在内核中：

__global__ void test(char** d_wordList) {
    printf("First character of the first word is: %c \n", d_wordList[0][0]);
}

内核应该得到一个指向每个单词开头的 int32 指针数组，实际上是一个 char**（或 int**），但它并没有像我预期的那样工作。

这种方法有什么问题？

一般来说，在 PyCUDA（甚至在 CUDA）中处理字符串的“标准”方法是什么？

提前致谢。

Answer 1

经过进一步思考，我得出结论，对于这种可变长度字符串的情况，在考虑数据问题时，使用“偏移数组”可能与二维索引（即双指针索引）没有太大区别在内核中访问。两者都涉及到一定程度的间接性。

这是一个演示这两种方法的工作示例：

$ cat t5.py
#!python
#!/usr/bin/env python
import time
import numpy as np
from pycuda import driver, compiler, gpuarray, tools
import math
from sys import getsizeof

import pycuda.autoinit

kernel_code1 = """
__global__ void test1(char** d_wordList) {
      (d_wordList[blockIdx.x][threadIdx.x])++;
}
    """

kernel_code2 = """
__global__ void test2(char* d_wordList, size_t *offsets) {
    (d_wordList[offsets[blockIdx.x] + threadIdx.x])++;
}
    """




mod = compiler.SourceModule(kernel_code1)
ker_test1 = mod.get_function("test1")



wordList = ['asd','bsd','csd']

d_words = []

for word in wordList:
    d_words.append(gpuarray.to_gpu(np.array(word, dtype=str)))

d_wordList = gpuarray.to_gpu(np.array([word.ptr for word in d_words], dtype=np.uintp))

ker_test1(d_wordList, block=(3,1,1), grid=(3,1,1))

for word in d_words:
  result = word.get()
  print result

mod2 = compiler.SourceModule(kernel_code2)
ker_test2 = mod2.get_function("test2")
wordlist2 = np.array(['asdbsdcsd'], dtype=str)
d_words2 = gpuarray.to_gpu(np.array(['asdbsdcsd'], dtype=str))
offsets = gpuarray.to_gpu(np.array([0,3,6,9], dtype=np.uint64))
ker_test2(d_words2, offsets, block=(3,1,1), grid=(3,1,1))
h_words2 = d_words2.get()
print h_words2


$ python t5.py
bte
cte
dte
['btectedte']
$

注意事项：

1. 对于双指针情况，与 OP 示例的唯一变化是使用 @talonmies 在 cmets 中建议的指针类型为 numpy.uintp

2. 我不认为数据的双指针访问一定会比与偏移查找方法相关的间接更快或更慢。另一项性能考虑是在将数据从主机复制到设备的领域，反之亦然。我相信双指针方法有效地涉及双向的多次分配和多次复制操作。对于很多字符串，这在主机/设备数据复制操作中会很明显。

3. 偏移方法的另一个可能的优点是它很容易确定每个字符串的长度 - 只需减去偏移列表中的两个相邻条目。这可能很有用，以便更容易确定有多少线程可以并行处理字符串，而不是让单个线程按顺序处理字符串（或使用内核代码中的方法来确定字符串长度，或通过每个字符串的长度）。