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

Posted 2021-04-21

我正在尝试使用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 
", d_wordList[0][0]);
}

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

这种方法有什么问题？

另外，在PyCUDA（甚至是CUDA）中使用字符串的“标准”方法是什么？

提前致谢。

答案

经过一番深思，我得出结论，对于这种可变长度字符串的情况，使用“偏移数组”可能与2D索引（即双指针索引）没有太大的不同，在考虑数据访问的问题时核心。两者都涉及一定程度的间接。

这是一个展示两种方法的实例：

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