可以将一系列无符号字符读为长时间，然后将其保存到另一组无符号字符中？

Question

input

输入，其中

unsigned char

是一个内核参数。 Nsight Compute报告说，内存访问不是合并的，在将每个字符从输入中保存到线程的本地阵列时会导致重大延迟。我的想法是一次读取数组8字节的块（作为一个

n * 57

），然后将它们存储在另一个64个字符的数组中，并与32位银行边界正确对齐。理想情况下，我应该使用共享内存来加快访问。 最初的内核是：n尽管我想删除填充物，而是像以前建议的那样使用共享内存：long

我认为可以正确实现，因为即使这不是正确的方法，第一个字符串也正确地读取了。

P.S

__global__ void sha256 (const unsigned char * __restrict__ input, unsigned char * __restrict__ hash, int n){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n)
        return;
    
    int len = device_strlen(input + idx * (MAX_WORD_LENGTH+1)); // Lunghezza del messaggio per il thread
    if (len > 56) {
        return;
    }

    unsigned int abh[8];
    unsigned int W[64];
    unsigned int temp1,temp2;
    unsigned int X[8]; 
    unsigned char padding[64];
    int i;

    //unrolling dei cicli per ottimizzare il codice evitare speculazione ecc
    #pragma unroll
    for(i = 0; i < 8 ; i++ )
        abh[i] = H[i];
    
        #pragma unroll
    for(i = 0; i < len; i++){
        padding[i] = input[idx * (MAX_WORD_LENGTH+1) + i];
    }

是64，

__global__ void sha256 (const unsigned char * __restrict__ input, unsigned char * __restrict__ hash, int n){
    // Identificatore del thread
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n)
        return;
    
    unsigned int abh[8];
    unsigned int temp1, temp2;
    unsigned long long dimensione;
    unsigned int W[64];

    __shared__ unsigned char shared_input[THREAD_PER_BLOCK * COALESCENT_WIDTH];
    int i;
    #pragma unroll
    for ( i = 0; i < COALESCENT_WIDTH / 8; i++) {
        *((unsigned long*)(shared_input + threadIdx.x * COALESCENT_WIDTH + i * 8)) = *((unsigned int*)(input + idx * (COALESCENT_WIDTH) + i * 8));
    }
    __syncthreads();
    // Calcolo della lunghezza della stringa
    int len = device_strlen(&shared_input[threadIdx.x*COALESCENT_WIDTH]);

是56 我必须在第二个内核输入中重新分配为

COALESCENT_WIDTH

MAX_WORD_LENGTH

的数组，才能像共享内存数组一样合并。
    

profiler报告没有合并内存，因为相邻的线程没有从内存中读取相邻单词。
您有以下行：

n * 64

LLET的邻近线程T0和T1。
让我们进行循环的第一次迭代（对于T0和T1）。
对于T0

char

，对于T1

Answer 1

。 从地址读取的两个线程

i=0

对于T0，它等于

idx=0

对于T1，它等于

idx=1

由于L1 CACHE为128字节，L2缓存为32个字节，从32个线程的翘曲中只有2个线程将从同一L1缓存线中读取，并且每个线程都会从单独的L2缓存线中读取。如果块中的其他线程尚未将其丢弃，则缓存将有助于随后的迭代步骤。 （与您的线程数量相比，缓存实际上很小）。

您想要的是在同一时刻阅读相邻单词。回到那天，我使用了这样的东西：

input + idx * (COALESCENT_WIDTH) + i * 8

您需要将其调整为

input + 0

类型，但是您明白了。