Cuda текстура выборки с пользовательским типом - PullRequest
Новая
       32

Cuda текстура выборки с пользовательским типом

1 голос
/ 08 мая 2019

Я хочу читать из памяти текстур, используя пользовательскую структуру (ushort8), я хочу читать 128 бит из каждой выборки.Когда я компилирую код, я получаю следующую ошибку: 

/usr/local/cuda-9.0/bin/nvcc -ccbin g++ -I /home/nvidia/NVIDIA_CUDA-9.0_Samples/common/inc/  -m64    --default-stream per-thread -Xptxas -v --resource-usage --maxrregcount=32 -gencode arch=compute_62,code=sm_62 -gencode arch=compute_62,code=compute_62 -o teste.o -c teste.cu #-DNO_COMPUTE_LOCAL
teste.cu(20): error: no instance of overloaded function "tex2DLayered" matches the argument list
            argument types are: (texture<ushort8, 242, cudaReadModeElementType>, const unsigned int, const unsigned int, const unsigned int)

Я использую cuda 9.0.

У меня уже есть версия, использующая глобальную память gpu, и я хочу повторить то же самое, используяТекстурная память.

Я уже пробовал с ushort4 и отлично работает.Вот код 

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>


typedef struct __align__(16) {
    unsigned short s0, s1, s2, s3, s4, s5, s6, s7;
}ushort8;


texture<ushort8, cudaTextureType2DLayered> d_samples;

__global__ void teste(){
    __shared__ ushort8 samples[4];

    samples[threadIdx.x]=tex2DLayered(d_samples,threadIdx.x,threadIdx.y,threadIdx.z);
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s0 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s1);
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s2 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s3 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s4 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s5 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s6 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s7 );

}


int main(int argc, char *argv[]){
    dim3 threadsPerBlock(4,1,1);
    dim3 numBlocks(1,1,1);
    cudaError_t err=cudaSuccess;
    cudaChannelFormatDesc channelDesc_samples = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindUnsigned);
    cudaArray *samples_3darray;

    ushort8 samples[4];

    samples[0].s0=0;
    samples[0].s1=1;
    samples[0].s2=2;
    samples[0].s3=3;
    samples[0].s4=4;
    samples[0].s5=5;
    samples[0].s6=6;
    samples[0].s7=7;

    samples[1].s0=8;
    samples[1].s1=9;
    samples[1].s2=10;
    samples[1].s3=11;
    samples[1].s4=12;
    samples[1].s5=13;
    samples[1].s6=14;
    samples[1].s7=15;

    samples[2].s0=16;
    samples[2].s1=17;
    samples[2].s2=18;
    samples[2].s3=19;
    samples[2].s4=20;
    samples[2].s5=21;
    samples[2].s6=22;
    samples[2].s7=23;

    samples[3].s0=24;
    samples[3].s1=25;
    samples[3].s2=26;
    samples[3].s3=27;
    samples[3].s4=28;
    samples[3].s5=29;
    samples[3].s6=30;
    samples[3].s7=31;


    err=cudaMalloc3DArray(&samples_3darray, &channelDesc_samples, make_cudaExtent(4, 1, 1), cudaArrayLayered);
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to allocate the texture memory for the samples(error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }
    cudaMemcpy3DParms myparms_samples = {0};
    myparms_samples.srcPos = make_cudaPos(0,0,0);
    myparms_samples.dstPos = make_cudaPos(0,0,0);
    myparms_samples.srcPtr = make_cudaPitchedPtr(samples, 4 * sizeof(ushort8),4, 1);
    myparms_samples.dstArray = samples_3darray;
    myparms_samples.extent = make_cudaExtent(4, 1,1);
    myparms_samples.kind = cudaMemcpyHostToDevice;

    d_samples.addressMode[0] = cudaAddressModeBorder;
    d_samples.addressMode[1] = cudaAddressModeBorder;
    d_samples.addressMode[2] = cudaAddressModeBorder;
    d_samples.filterMode = cudaFilterModePoint;
    d_samples.normalized = false;  // access with normalized texture coordinates
    err=cudaMemcpy3D(&myparms_samples);
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to copy the image samples from host to device (error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }

    err=cudaBindTextureToArray(d_samples, samples_3darray, channelDesc_samples);
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to bind the texture memory (error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }

    teste<<<numBlocks, threadsPerBlock>>>();
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to launch the kernel for the calculation of the local sum (error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }

    err=cudaUnbindTexture(d_samples);
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to unbind the image(error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }
    err=cudaFreeArray(samples_3darray);
    if(err!=cudaSuccess){
        fprintf(stderr, "Failed to free the samples(error code %d)!\n", cudaGetLastError());
        exit(EXIT_FAILURE);
    }

    return 0;

}

Кто-нибудь может мне помочь?спасибо!

1 Ответ

2 голосов
/ 08 мая 2019

Я думаю, я нашел обходной путь.Я использовал reinterpret_cast, и он добился цели. 

texture<uint4, cudaTextureType2DLayered> d_samples;

__global__ void teste(){
    __shared__ ushort8 samples[4];

    reinterpret_cast<uint4*>(samples)[threadIdx.x]=tex2DLayered(d_samples,threadIdx.x,threadIdx.y,threadIdx.z);
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s0 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s1);
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s2 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s3 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s4 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s5 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s6 );
    printf("thread=%d, value=%hu\n",threadIdx.x, samples[threadIdx.x].s7 );

}
...