wordpress   发布时间:2022-04-02  发布网站:大佬教程  code.js-code.com
大佬教程收集整理的这篇文章主要介绍了如何在CUDA中使用uint4向量正确转换全局内存数组以增加内存吞吐量?大佬教程大佬觉得挺不错的,现在分享给大家,也给大家做个参考。

概述

通常有两种技术可以在计算能力1.3 GPU上增加CUDA内核上全局内存的内存吞吐量.内存访问合并并访问至少4个字节的字.利用第一种技术,通过相同半warp的线程对相同存储器段的访问被合并为更少的事务,而在访问至少4字节的字时,该存储器段有效地从32字节增加到128. 更新:基于talonmies回答的解决方案.当存在全局存储器中存在无符号字符时,要访问16字节而不是1字节字,通常通过将存储器阵列转
通常有两种技术可以在计算能力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("%c\n",h_out[i]);

    return 0;
}@H_673_24@

解决方法

如果我已经理解了你要做的事情,那么逻辑方法是使用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 %c\n",i,m[i]);

    return 0;
}@H_673_24@ 
 

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

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

大佬总结

以上是大佬教程为你收集整理的如何在CUDA中使用uint4向量正确转换全局内存数组以增加内存吞吐量?全部内容,希望文章能够帮你解决如何在CUDA中使用uint4向量正确转换全局内存数组以增加内存吞吐量?所遇到的程序开发问题。

如果觉得大佬教程网站内容还不错,欢迎将大佬教程推荐给程序员好友。

本图文内容来源于网友网络收集整理提供,作为学习参考使用,版权属于原作者。
如您有任何意见或建议可联系处理。小编QQ:384754419,请注明来意。