CUDA 配列に格納する必要がある次の文字列があるとします。
"やあ"
"これは"
"誰が"
これを行うには、GPU で配列を宣言する方法を教えてください。C++ を使ってみstrings
ましたが、うまくいきません。
おそらくこれを行う最善の方法は、一般的な圧縮された疎行列形式に似た構造を使用することです。文字データを 1 つの線形メモリにパックして格納し、別の整数配列を使用して開始インデックスを格納し、おそらく 3 番目の配列を使用して文字列の長さを格納します。後者のストレージ オーバーヘッドは、データ内のすべてのエントリに対して文字列終了バイトを格納し、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
。Google で「CUDA string」を検索すると、この CUDA の「Hello World」の例が最初のヒットとして
表示されます: http://computer-graphics.se/hello-world-for-cuda.htmlchar*
CUDA の文字列。CUDA では標準の C 関数が使用できないことに注意してくださいstrcpy
。strcmp
文字列の配列が必要な場合は、char**
(C/C++ のように) を使用するだけです。および同様の機能についてstrcmp
は、何をしたいかによって大きく異なります。CUDA は文字列操作にはあまり適していません。何をしたいのかについてもう少し詳しく教えていただけると助かります。