배열을 텍스처 표현으로 변환

Question

나는 cuda ....에서 tex2D로 배열을 나타내려고 노력 중이다. 디버깅을 한 후, 1 백만 개의 요소 중 19 개가 텍스처에 잘못 복사되었다는 것을 알았다. 배열, 나는 0 intstead 1 있어요. 난 당신의 코드에서 볼 수

void evolve_gpu(byte* h_in, byte* h_out) 
{ 

//int SIZE = N * N * N * N * sizeof(float); 
cudaEvent_t start, stop; 
size_t d_in_pitch; 
size_t d_out_pitch; 
int len = 1002; 

checkCudaErrors(cudaEventCreate(&start)); 
checkCudaErrors(cudaEventCreate(&stop)); 

// Allocate the device input image array 
unsigned char *d_in = NULL; 
unsigned char *d_out = NULL; 
checkCudaErrors(cudaMallocPitch(&d_in, &d_in_pitch, sizeof(unsigned char)*len, len)); 
checkCudaErrors(cudaMallocPitch(&d_out, &d_out_pitch, sizeof(unsigned char)*len, len)); 

// Copy the host input image to the device memory 
checkCudaErrors(cudaMemcpy2D(d_in, d_in_pitch, h_in, sizeof(unsigned char)*len 
    , sizeof(unsigned char)*len, len, cudaMemcpyHostToDevice)); 




/**************************** TEXTURE CONFIGURATION ******************************/ 
cudaResourceDesc resDesc; 
memset(&resDesc, 0, sizeof(resDesc)); 
resDesc.resType = cudaResourceTypePitch2D; 
resDesc.res.pitch2D.devPtr = d_in; 
resDesc.res.pitch2D.pitchInBytes = d_in_pitch; 
resDesc.res.pitch2D.width = len; 
resDesc.res.pitch2D.height = len; 
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<unsigned char>(); 

cudaTextureDesc texDesc; 
memset(&texDesc, 0, sizeof(texDesc)); 
texDesc.readMode = cudaReadModeElementType; 
texDesc.normalizedCoords=false; 
texDesc.addressMode[0]=cudaAddressModeBorder; 
texDesc.addressMode[1]=cudaAddressModeBorder; 

cudaTextureObject_t tex; 
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL); 
/*********************************************************************************/ 

checkCudaErrors(cudaEventRecord(start, NULL)); 

// Launch the CUDA Kernel 
dim3 block = dim3(THREADS_X, THREADS_Y); 
dim3 grid = dim3((len+block.x-1)/block.x,(len+block.y-1)/block.y);//25*50 
evolve_kernel<<<grid, block>>>(tex, d_out); 

//******** kernel<<< number of blocks, number of threads, dynamic memory per block, associated stream >>> *******// 

// Copy the device result to the host 
checkCudaErrors(cudaMemcpy2D(h_out, d_out_pitch, 
          d_out, d_out_pitch, 
          sizeof(unsigned char)*len, len, 
          cudaMemcpyDeviceToHost));  
for(int i=0;i<1002*1002;i++){ 

    if(h_in[i] != h_out[i]) 
     printf("i = %d\n",i); 


} 
checkCudaErrors(cudaGetLastError()); 

checkCudaErrors(cudaEventRecord(stop, NULL)); 
checkCudaErrors(cudaEventSynchronize(stop)); 

checkCudaErrors(cudaFree(d_in)); 
checkCudaErrors(cudaFree(d_out)); 

float msec = 0.f; 
checkCudaErrors(cudaEventElapsedTime(&msec, start, stop)); 

printf("Basic version took: %f ms\n", msec); 

} 

Answer 1

하나 개에 문제가 장치 -> 호스트 카피 :

checkCudaErrors(cudaMemcpy2D(h_out, d_out_pitch, 
         d_out, d_out_pitch, 
         sizeof(unsigned char)*len, len, 
         cudaMemcpyDeviceToHost)); 

documentation을 참조하면,이 cudaMemcpy2D 호출을하기위한 곳 중 2 번째 매개 변수는 대상의 피치입니다 할당 (즉, 피치가 h_out 인 경우). 그러나 h_out은 피치 할당을 의미하지는 않을 것이며, 어쨌든 그 피치는 d_out_pitch에 의해 주어지지 않을 것입니다.

checkCudaErrors(cudaMemcpy2D(h_out, len*sizeof(unsigned char), 
         d_out, d_out_pitch, 
         sizeof(unsigned char)*len, len, 
         cudaMemcpyDeviceToHost)); 

I :

당신은 h_out 및 h_in는, 그 두번째 매개 변수는 h_out 배열의 (UN 피치의) 폭을 변경해야 유사한 할당은 가정, 완전한 코드를 표시하지 않았지만

는

evolve_kernel<<<grid, block>>>(tex, d_out); 

나는 전자를 할 것이다 : '또 커널은 당신이 그것에 d_out의 피치를 통과하지 않는 d_out (A 투구 할당)에서 제대로 작동 할 수 있는지 궁금 해요

evolve_kernel<<<grid, block>>>(tex, d_out, d_out_pitch); 

으로 전화를 걸었지만 커널 코드가 표시되지 않았습니다. 것 같다

$ cat t648.cu 
#include <stdio.h> 
#include <helper_cuda.h> 
#define THREADS_X 16 
#define THREADS_Y 16 

const int len = 1002; 
typedef unsigned char byte; 

__global__ void evolve_kernel(cudaTextureObject_t tex, unsigned char *d_out, size_t pitch){ 
    int idx = threadIdx.x+blockDim.x*blockIdx.x; 
    int idy = threadIdx.y+blockDim.y*blockIdx.y; 
    if ((idx < len) && (idy < len)) 
    d_out[idy*pitch+idx] = tex2D<unsigned char>(tex, idx, idy); 
} 



void evolve_gpu(byte* h_in, byte* h_out) 
{ 

//int SIZE = N * N * N * N * sizeof(float); 
    cudaEvent_t start, stop; 
    size_t d_in_pitch; 
    size_t d_out_pitch; 

    checkCudaErrors(cudaEventCreate(&start)); 
    checkCudaErrors(cudaEventCreate(&stop)); 

// Allocate the device input image array 
    unsigned char *d_in = NULL; 
    unsigned char *d_out = NULL; 
    checkCudaErrors(cudaMallocPitch(&d_in, &d_in_pitch, sizeof(unsigned char)*len, len)); 
    checkCudaErrors(cudaMallocPitch(&d_out, &d_out_pitch, sizeof(unsigned char)*len, len)); 

// Copy the host input image to the device memory 
    checkCudaErrors(cudaMemcpy2D(d_in, d_in_pitch, h_in, sizeof(unsigned char)*len 
    , sizeof(unsigned char)*len, len, cudaMemcpyHostToDevice)); 




/**************************** TEXTURE CONFIGURATION ******************************/ 
    cudaResourceDesc resDesc; 
    memset(&resDesc, 0, sizeof(resDesc)); 
    resDesc.resType = cudaResourceTypePitch2D; 
    resDesc.res.pitch2D.devPtr = d_in; 
    resDesc.res.pitch2D.pitchInBytes = d_in_pitch; 
    resDesc.res.pitch2D.width = len; 
    resDesc.res.pitch2D.height = len; 
    resDesc.res.pitch2D.desc = cudaCreateChannelDesc<unsigned char>(); 

    cudaTextureDesc texDesc; 
    memset(&texDesc, 0, sizeof(texDesc)); 
    texDesc.readMode = cudaReadModeElementType; 
    texDesc.normalizedCoords=false; 
    texDesc.addressMode[0]=cudaAddressModeBorder; 
    texDesc.addressMode[1]=cudaAddressModeBorder; 

    cudaTextureObject_t tex; 
    cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL); 
/*********************************************************************************/ 

    checkCudaErrors(cudaEventRecord(start, NULL)); 

// Launch the CUDA Kernel 
    dim3 block = dim3(THREADS_X, THREADS_Y); 
    dim3 grid = dim3((len+block.x-1)/block.x,(len+block.y-1)/block.y);//25*50 
    evolve_kernel<<<grid, block>>>(tex, d_out, d_out_pitch); 

//******** kernel<<< number of blocks, number of threads, dynamic memory per block, associated stream >>> *******// 

// Copy the device result to the host 
    checkCudaErrors(cudaMemcpy2D(h_out, len*sizeof(unsigned char), 
          d_out, d_out_pitch, 
          sizeof(unsigned char)*len, len, 
          cudaMemcpyDeviceToHost)); 
    for(int i=0;i<1002*1002;i++){ 

    if(h_in[i] != h_out[i]) 
     printf("i = %d\n",i); 


    } 
    checkCudaErrors(cudaGetLastError()); 

    checkCudaErrors(cudaEventRecord(stop, NULL)); 
    checkCudaErrors(cudaEventSynchronize(stop)); 

    checkCudaErrors(cudaFree(d_in)); 
    checkCudaErrors(cudaFree(d_out)); 

    float msec = 0.f; 
    checkCudaErrors(cudaEventElapsedTime(&msec, start, stop)); 

    printf("Basic version took: %f ms\n", msec); 

} 

int main(){ 
    byte *h_data_in, *h_data_out; 
    h_data_in = (byte *)malloc(len*len*sizeof(byte)); 
    h_data_out = (byte *)malloc(len*len*sizeof(byte)); 
    for (int i = 0; i < len*len; i++){ 
    h_data_in[i] = 3; 
    h_data_out[i] = 0;} 
    evolve_gpu(h_data_in, h_data_out); 
    return 0; 
} 
$ nvcc -arch=sm_35 -I/usr/local/cuda/samples/common/inc t648.cu -o t648 
$ ./t648 
Basic version took: 3.868576 ms 
$ 

이 제대로 작동 시험을 통과 :

는 여기에 내가 예를 구축하는 고정 위의 문제와 몇 가지 다른 변화는 보여 코드, 주위에 만들어 완벽하게 일했다 예제 당신이 만들었습니다. 이 유형의 질문 ("이 코드가 작동하지 않는 이유는 무엇입니까?")에 대해 [MCVE] (http : //stackoverflow.com/help/on-topic)를 제공하려면