PyCUDA 二维数组实现（或使用字符串）

Question

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

我尝试通过创建单个字符串来展平它，但随后要为它编制索引，每次内核运行时我都必须检查其中的一些内容。如果有 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**），但它没有像我预期的那样工作。

这种方法有什么问题？

还有 "standard" 在 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']
$

备注：

对于双指针的情况，OP 示例的唯一变化是使用 numpy.uintp 类型作为指针，正如@talonmies
[在评论中所建议的那样]
我认为数据的双指针访问不一定比与偏移查找方法关联的间接访问更快或更慢。另一个性能考虑因素是将数据从主机复制到设备，反之亦然。我相信，双指针方法在两个方向上有效地涉及多个分配和多个复制操作。对于很多字符串，这在 host/device 数据复制操作中会很明显。
偏移方法的另一个可能优点是很容易确定每个字符串的长度——只需减去偏移列表中的两个相邻条目即可。这可能很有用，可以很容易地确定有多少线程可以并行操作一个字符串，而不是让单个线程顺序地处理一个字符串（或使用内核代码中的方法来确定字符串长度，或通过每个字符串的长度）。

PyCUDA 二维数组实现（或使用字符串）

PyCUDA 2D array implementations (or working with strings)

python

arrays

string

cuda

pycuda