CUDA已經built-in variables像threadIdx.x
和blockIdx.x
識別塊內的每個線程,也阻斷該線程是一種塊內的每個線程都將具有相同的blockIdx.x
值,但塊中的每個線程將具有不同的/唯一的(每個塊)threadIdx.x
值。
因此,我們可以使用blockIdx.x
爲每個塊選擇一個特定的序列。這個變量可以用來選擇正確的序列長度以及每個序列/塊的偏移量。
我們可以爲每個序列項/字符分配一個線程。我們可以使用threadIdx.x
來標識每個線程應該選擇哪個序列成員。
這裏是一個完全樣例:
$ cat t405.cu
#include <stdio.h>
__global__ void tk(char *seq, int *offsets, int *seq_lengths, int num_seq){
if (blockIdx.x < num_seq)
if (threadIdx.x < seq_lengths[blockIdx.x])
printf("block: %d, thread: %d, seq: %c\n", blockIdx.x, threadIdx.x, seq[offsets[blockIdx.x]+threadIdx.x]);
}
int main(){
char seq[] = {'a','b','f','g','c','d','>','b','g','d','>','a','b', 'c', 'd', 'e', '>'};
int seq_length[] = { 6, 3, 5 };
int offsets[] = { 0, 7, 11 };
int num_seq = 3;
int seq_sz = sizeof(seq);
int seq_l_sz = sizeof(seq_length);
int off_sz = sizeof(offsets);
char *d_seq;
int *d_seq_length, *d_offsets;
cudaMalloc(&d_seq, seq_sz);
cudaMalloc(&d_seq_length, seq_l_sz);
cudaMalloc(&d_offsets, off_sz);
cudaMemcpy(d_seq, seq, seq_sz, cudaMemcpyHostToDevice);
cudaMemcpy(d_seq_length, seq_length, seq_l_sz, cudaMemcpyHostToDevice);
cudaMemcpy(d_offsets, offsets, off_sz, cudaMemcpyHostToDevice);
tk<<<num_seq, 1024>>>(d_seq, d_offsets, d_seq_length, num_seq);
cudaDeviceSynchronize();
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err) printf("cuda error: %s\n", cudaGetErrorString(err));
return 0;
}
$ nvcc -arch=sm_61 -o t405 t405.cu
$ ./t405
block: 1, thread: 0, seq: b
block: 1, thread: 1, seq: g
block: 1, thread: 2, seq: d
block: 2, thread: 0, seq: a
block: 2, thread: 1, seq: b
block: 2, thread: 2, seq: c
block: 2, thread: 3, seq: d
block: 2, thread: 4, seq: e
block: 0, thread: 0, seq: a
block: 0, thread: 1, seq: b
block: 0, thread: 2, seq: f
block: 0, thread: 3, seq: g
block: 0, thread: 4, seq: c
block: 0, thread: 5, seq: d
$
如果您預計超過1024個字符的序列,那麼就要修改上面,也許是爲了讓每個線程處理多個角色,也許在循環。