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个字符的序列,那么就要修改上面,也许是为了让每个线程处理多个角色,也许在循环。