如何在CUDA中使用uint4向量正确转换全局内存数组以增加内存吞吐

发布时间：2020-12-14 04:37:23 所属栏目：Windows 来源：网络整理

导读：通常有两种技术可以在计算能力1.3 GPU上增加CUDA内核上全局内存的内存吞吐量.内存访问合并并访问至少4个字节的字.利用第一种技术,通过相同半warp的线程对相同存储器段的访问被合并为更少的事务,而在访问至少4字节的字时,该存储器段有效地从32字节增加到128.

通常有两种技术可以在计算能力1.3 GPU上增加CUDA内核上全局内存的内存吞吐量.内存访问合并并访问至少4个字节的字.利用第一种技术,通过相同半warp的线程对相同存储器段的访问被合并为更少的事务,而在访问至少4字节的字时,该存储器段有效地从32字节增加到128.

更新：基于talonmies回答的解决方案.当存在全局存储器中存在无符号字符时,要访问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,cudaMemcpyDeviceToHost );

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

    return 0;
}

解决方法

如果我已经理解了你要做的事情,那么逻辑方法是使用C reinterpret_cast机制使编译器生成正确的向量加载指令,然后使用内置字节大小的向量类型uchar4的CUDA来访问每个内部的每个字节.从全局内存加载的四个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 %cn",i,m[i]);

    return 0;
}

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

需要注意的是,用于计算1.x目标的open64编译器通常会失败这种尝试生成向量加载的策略,如果它能够检测到并非向量中的所有单词都被实际使用.因此,请确保触摸输入向量类型中的所有输入字以确保编译器可以正常播放.

（编辑：李大同）

【声明】本站内容均来自网络，其相关言论仅代表作者个人观点，不代表本站立场。若无意侵犯到您的权利，请及时与联系站长删除相关内容!