Le bloc lit une longueur de séquence différente gpu

Question

J'ai un tableau qui a une longueur de séquence différente, chaque séquence se terminant par '>'. seq = [a, b, f, g, c, d,>, b, g, d,> ....]. J'ai calculé la longueur de chaque séquence et l'ai stockée dans un tableau différent appelé seq_length = [6,3,5, ...]. Ensuite, j'utilise l'analyse exclusive pour calculer le décalage et l'enregistre dans un tableau appelé offset = [0, 6, 9, ..]. Ce que je veux, c'est que chaque bloc lise une séquence du tableau seq [] en utilisant la valeur de décalage. Par exemple, le bloc 0 lit la séquence qui commence à partir de seq [0] et s'arrête quand la longueur est = 6, le bloc 1 lit la séquence qui commence à partir de seq [6] et s'arrête quand la longueur est = 3, et ainsi de suite. Comment je peux le faire dans CUDA ?? Comment je peux laisser chaque bloc lire à partir d'une entrée de tableau différente (tableau seq []).

Merci d'avance :)

Answer 1

CUDA a built-in variables comme threadIdx.x et blockIdx.x qui identifient chaque fil dans un bloc, et aussi qui bloquent ce fil est. Chaque thread dans un bloc aura la même valeur blockIdx.x, mais chaque thread dans un bloc aura une valeur différente/unique (par bloc) threadIdx.x.

Par conséquent, nous pouvons utiliser blockIdx.x pour sélectionner une séquence spécifique pour chaque bloc. Cette variable peut être utilisée pour sélectionner la longueur de séquence correcte ainsi que le décalage pour chaque séquence/bloc.

Nous pouvons affecter un fil par bloc à chaque élément/caractère de séquence. Nous pouvons utiliser threadIdx.x pour identifier, pour chaque thread, quel membre de séquence il doit sélectionner.

Voici un exemple bien travaillé:

$ 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 
$ 

Si vous attendez plus de 1024 caractères dans une séquence, alors vous voulez modifier ce qui précède, peut-être d'avoir chaque thread gérer plusieurs caractères, peut-être dans un boucle.