4

通常有两种技术可以提高计算能力 1.3 GPU 上 CUDA 内核上全局内存的内存吞吐量；内存访问合并和访问至少 4 个字节的字。使用第一种技术，通过相同半扭曲的线程访问相同的内存段被合并为更少的事务，同时访问至少 4 个字节的字，这个内存段有效地从 32 个字节增加到 128 个。

更新：基于 talonmies answer 的解决方案。当全局内存中存储有无符号字符时，要访问 16 字节而不是 1 字节字，通常通过将内存数组转换为 uint4 来使用 uint4 向量。要从 uint4 向量中获取值，可以将其重铸为 uchar4，如下所示：

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>

__global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) {

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ unsigned char s_array[];

    uint4 *uint4_text = reinterpret_cast<uint4 *>(d_text);
    uint4 uint4_var;

    //memory transaction
    uint4_var = uint4_text[0];

    //recast data to uchar4
    uchar4 c0 = *reinterpret_cast<uchar4 *>(&uint4_var.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&uint4_var.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&uint4_var.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&uint4_var.w);

    d_out[idx] = c0.y;
}

int main ( void ) {

    unsigned char *d_text, *d_out;

    unsigned char *h_out = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) );
    unsigned char *h_text = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) );

    int i;

    for ( i = 0; i < 16; i++ )
            h_text[i] = 65 + i;

    cudaMalloc ( ( void** ) &d_text, 16 * sizeof ( unsigned char ) );
    cudaMalloc ( ( void** ) &d_out, 16 * sizeof ( unsigned char ) );

    cudaMemcpy ( d_text, h_text, 16 * sizeof ( unsigned char ), cudaMemcpyHostToDevice );

    kernel<<<1,16>>>(d_text, d_out );

    cudaMemcpy ( h_out, d_out, 16 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost );

    for ( i = 0; i < 16; i++ )
            printf("%c\n", h_out[i]);

    return 0;
}
4

2 回答 2

3

转换为 char* 就可以了。你试过了吗？如果是这样，是什么引发了这个问题？

在您的示例中，您似乎可以只s_array转换为 anint*并从中执行单个副本var.x（乘以j4 而不是 16）。

如果您需要更灵活地对单词中的字节进行改组，则可以使用__byte_perm()内在函数。例如，要反转整数中字节的顺序x，您可以这样做__byte_perm(x, 0, 0x0123);

使用向量类型甚至单个 int 来存储字节可能不会获得任何好处。在 Fermi 上，全局内存事务为 128 字节宽。因此，当您的 warp 遇到从/向全局内存加载/存储的指令时，GPU 将执行服务 32 个线程所需的 128 字节事务。性能很大程度上取决于需要多少单独的事务，而不是每个线程如何加载或存储其字节。

于 2012-10-27T18:43:35.290 回答
3

如果我了解您要做什么，逻辑方法是使用 C++reinterpret_cast机制使编译器生成正确的向量加载指令，然后使用 CUDA 内置字节大小的向量类型uchar4来访问四个字节中的每个字节从全局内存中加载的 32 位字。使用这种方法，您真的相信编译器知道在每个 32 位寄存器中进行字节访问的最佳方式。

一个完全人为的示例可能如下所示：

#include <cstdio>
#include <cstdlib>

__global__
void kernel(unsigned int *in, unsigned char* out)
{
    int tid = threadIdx.x;

    uint4* p = reinterpret_cast<uint4*>(in);
    uint4  i4 = p[tid]; // vector load here

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&i4.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&i4.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&i4.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&i4.w);

    out[tid*4+0] = c0.x;
    out[tid*4+1] = c4.y;
    out[tid*4+2] = c8.z;
    out[tid*4+3] = c12.w;
}

int main(void)
{
    unsigned int c[8] = { 
        2021161062, 2021158776, 2020964472, 1920497784, 
        2021161058, 2021161336, 2020898936, 1702393976 };

    unsigned int * _c;
    cudaMalloc((void **)&_c, sizeof(int)*size_t(8));
    cudaMemcpy(_c, c, sizeof(int)*size_t(8), cudaMemcpyHostToDevice);
    unsigned char * _m;
    cudaMalloc((void **)&_m, sizeof(unsigned char)*size_t(8));

    kernel<<<1,2>>>(_c, _m);

    unsigned char m[8];
    cudaMemcpy(m, _m, sizeof(unsigned char)*size_t(8), cudaMemcpyDeviceToHost);

    for(int i=0; i<8; i++)
        fprintf(stdout, "%d %c\n", i, m[i]);

    return 0;
}

它应该产生一个嵌入提供给内核的无符号整数数组中的可读字符串。

一个警告是，用于计算 1.x 目标的 open64 编译器通常会破坏这种尝试生成向量负载的策略，如果它可以检测到并非向量中的所有单词都被实际使用。因此，请确保触摸输入向量类型中的所有输入词，以确保编译器运行良好。

于 2012-10-28T12:27:06.780 回答