C/CUDA：只能索引 CudaArray 中的每四个元素

Question

这是我的第一篇文章，所以我很高兴能获得一些新的见解并扩大我的知识。目前我正在开发一个 C 项目，其中加载了一个带有 3d 数据的二进制原始文件，在 CUDA 中处理并保存在一个新的二进制原始文件中。

这是基于 CUDA Samples 中的 simpleTexture3D 项目：

这是我的cpp

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

// includes, cuda
#include <vector_types.h>
#include <driver_functions.h>
#include <cuda_runtime.h>

// CUDA utilities and system includes
#include <helper_cuda.h>
#include <helper_functions.h>
#include <vector_types.h>

typedef unsigned int  uint;
typedef unsigned char uchar;

const char *sSDKsample = "simpleTexture3D";

const char *volumeFilename = "Bucky.raw";
const cudaExtent volumeSize = make_cudaExtent(32, 32, 32);

const uint width = 64, height = 64, depth=64;

//const char *volumeFilename = "TestOCT.raw";
//const cudaExtent volumeSize = make_cudaExtent(1024, 512, 512);
//
//const uint width = 1024, height = 512, depth=512;
const dim3 blockSize(8, 8, 8);
const dim3 gridSize(width / blockSize.x, height / blockSize.y, depth / blockSize.z);

uint *d_output = NULL;

int *pArgc = NULL;
char **pArgv = NULL;

extern "C" void cleanup();
extern "C" void initCuda(const uchar *h_volume, cudaExtent volumeSize);
extern "C" void render_kernel(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH, uint imageD);

void loadVolumeData(char *exec_path);


// render image using CUDA
void render()
{
    // call CUDA kernel
    render_kernel(gridSize, blockSize, d_output, width, height, depth);

    getLastCudaError("render_kernel failed");
}

void cleanup()
{
    // cudaDeviceReset causes the driver to clean up all state. While
    // not mandatory in normal operation, it is good practice.  It is also
    // needed to ensure correct operation when the application is being
    // profiled. Calling cudaDeviceReset causes all profile data to be
    // flushed before the application exits
    checkCudaErrors(cudaDeviceReset());
}

// Load raw data from disk
uchar *loadRawFile(const char *filename, size_t size)
{
    FILE *fp = fopen(filename, "rb");

    if (!fp)
    {
        fprintf(stderr, "Error opening file '%s'\n", filename);
        return 0;
    }

    uchar *data = (uchar *) malloc(size);
    size_t read = fread(data, 1, size, fp);
    fclose(fp);

    printf("Read '%s', %lu bytes\n", filename, read);

    return data;
}

// write raw data to disk
int writeRawFile(const char *filename, uchar *data, size_t size)
{
    int returnState=0;
    // cut file extension from filename
    char *a=strdup(filename);   //via strdup you dumb a const char to char, you must free it yourself

    int len = strlen(a);
    a[len-4] = '\0';    //deletes '.raw'
    //printf("%s\n",a);

    char b[50];
    sprintf(b, "_%dx%dx%d_out.raw", width, height, depth);

    //char b[]="_out.raw";  //Add suffix out to filename

    char buffer[256]; // <- danger, only storage for 256 characters.
    strncpy(buffer, a, sizeof(buffer));
    strncat(buffer, b, sizeof(buffer));
    free(a);

    FILE *fp = fopen(buffer, "wb"); //Open or create file for writing as binary, all existing data is cleared

    if (!fp)
    {
        fprintf(stderr, "Error opening or creating file '%s'\n", buffer);
        return 0;
    }

    size_t write = fwrite(data, 1, size, fp);
    fclose(fp);

    if (write==size)
    {    
         printf("Wrote %lu bytes to '%s'\n", write, buffer);
         return 0;
    }
    else
    {
        printf("Error writing data to file '%s'\n", buffer);    
        return 1;
    }
}

// General initialization call for CUDA Device
int chooseCudaDevice(int argc, char **argv)
{
    int result = 0;

    result = findCudaDevice(argc, (const char **)argv);

    return result;
}

void runAutoTest(char *exec_path, char *PathToFile)
{
    // set path
    char *path;
    if (PathToFile == NULL)
    {
        path = sdkFindFilePath(volumeFilename, exec_path);
    }
    else
    {
        path = PathToFile;
    }

    if (path == NULL)
    {
        fprintf(stderr, "Error unable to find 3D Volume file: '%s'\n", volumeFilename);
        exit(EXIT_FAILURE);
    }

    // Allocate output memory
    checkCudaErrors(cudaMalloc((void **)&d_output, width*height*depth*sizeof(uchar)));

    // zero out the output array with cudaMemset
    cudaMemset(d_output, 0, width*height*depth*sizeof(uchar));

    // render the volumeData
    render_kernel(gridSize, blockSize, d_output, width, height, depth);

    checkCudaErrors(cudaDeviceSynchronize());
    getLastCudaError("render_kernel failed");

    uchar *h_output = (uchar*)malloc(width*height*depth);
    checkCudaErrors(cudaMemcpy(h_output, d_output, width*height*depth*sizeof(uchar), cudaMemcpyDeviceToHost));
    int wState=writeRawFile(path,h_output,width*height*depth);

    checkCudaErrors(cudaFree(d_output));
    free(h_output);

    // cudaDeviceReset causes the driver to clean up all state. While
    // not mandatory in normal operation, it is good practice.  It is also
    // needed to ensure correct operation when the application is being
    // profiled. Calling cudaDeviceReset causes all profile data to be
    // flushed before the application exits
    cudaDeviceReset();

    //exit(bTestResult ? EXIT_SUCCESS : EXIT_FAILURE);
}


void loadVolumeData(char *exec_path, char *PathToFile)
{
    char *path;
    // load volume data
    if (PathToFile == NULL)
    {
        path = sdkFindFilePath(volumeFilename, exec_path);
    }
    else
    {
        path = PathToFile;
    }

    if (path == NULL)
    {
        fprintf(stderr, "Error unable to find 3D Volume file: '%s'\n", volumeFilename);
        exit(EXIT_FAILURE);
    }

    size_t size = volumeSize.width*volumeSize.height*volumeSize.depth;
    uchar *h_volume = loadRawFile(path, size);
    //int wState=writeRawFile(path,h_volume,size);

    initCuda(h_volume, volumeSize);

    free(h_volume);
}


////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main(int argc, char **argv)
{
    pArgc = &argc;
    pArgv = argv;

    char *image_file = NULL;

    printf("%s Starting...\n\n", sSDKsample);

    if (checkCmdLineFlag(argc, (const char **)argv, "file"))    //Note cmd line argument is -file "PathToFile/File.raw"
    {                                                           // for example -file "C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.0\2_Graphics\simpleTexture3D_FanBeamCorr\data\TestOCT_Kopie.raw"            
        getCmdLineArgumentString(argc, (const char **)argv, "file", &image_file);
    }

    if (image_file)
    {
        chooseCudaDevice(argc, argv);

        loadVolumeData(argv[0],image_file);

        runAutoTest(argv[0],image_file);
    }
    else
    {
        // use command-line specified CUDA device, otherwise use device with highest Gflops/s
        chooseCudaDevice(argc, argv);

        loadVolumeData(argv[0],NULL);

        runAutoTest(argv[0],NULL);
    }

    printf("I am finished...\n"
           "Can I get some ice cream please\n");

    exit(EXIT_SUCCESS);
}

这是我的 .cu

#ifndef _SIMPLETEXTURE3D_KERNEL_CU_
#define _SIMPLETEXTURE3D_KERNEL_CU_


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

#include <helper_cuda.h>
#include <helper_math.h>

typedef unsigned int  uint;
typedef unsigned char uchar;

texture<uchar, 3, cudaReadModeNormalizedFloat> tex;  // 3D texture

cudaArray *d_volumeArray = 0;

__global__ void
d_render(uint *d_output, uint imageW, uint imageH, uint imageD)
{
    uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
    uint z = __umul24(blockIdx.z, blockDim.z) + threadIdx.z;

 //   float u = x / (float) imageW;
 //   float v = y / (float) imageH;
    //float w = z / (float) imageD;
 //   // read from 3D texture
 //   float voxel = tex3D(tex, u, v, w);
    uint ps=__umul24(imageW,imageH);

    if ((x < imageW) && (y < imageH) && (z < imageD))
    {
        // write output color
        uint i = __umul24(z,ps) +__umul24(y, imageW) + x;
        d_output[1] = (uchar) 255;//+0*voxel*255;
    }
}

extern "C"
void initCuda(const uchar *h_volume, cudaExtent volumeSize)
{
    // create 3D array
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
    checkCudaErrors(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize));

    // copy data to 3D array
    cudaMemcpy3DParms copyParams = {0};
    copyParams.srcPtr   = make_cudaPitchedPtr((void *)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
    copyParams.dstArray = d_volumeArray;
    copyParams.extent   = volumeSize;
    copyParams.kind     = cudaMemcpyHostToDevice;
    checkCudaErrors(cudaMemcpy3D(&copyParams));

    // set texture parameters
    tex.normalized = true;                      // access with normalized texture coordinates
    tex.filterMode = cudaFilterModeLinear;      // linear interpolation
    tex.addressMode[0] = cudaAddressModeBorder;   // wrap texture coordinates
    tex.addressMode[1] = cudaAddressModeBorder;
    tex.addressMode[2] = cudaAddressModeBorder;

    // bind array to 3D texture
    checkCudaErrors(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));
}

extern "C"
void render_kernel(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH, uint imageD)
{
    d_render<<<gridSize, blockSize>>>(d_output, imageW, imageH, imageD);
}

#endif // #ifndef _SIMPLETEXTURE3D_KERNEL_CU_

如您所见，目前，除了 index = 1（设置为 255）之外，我将所有值都设置为零。然而，当我现在在斐济打开图像堆栈时，我看到第一张幻灯片上的第四个像素是白色的。如果我改用 index=i，我会每隔四列周期性地在图像堆栈上获得白色垂直线。一般来说，似乎只有每四个元素在 CudaArray 中被索引。所以我想知道这里是否存在由 sizeof(uchar)=1 和 sizeof(uint)=4 导致的某种错误。显然会有因素 4 :)

我渴望从你们这里得到专家

干杯米卡

Answer 1

我自己想出来的。内核与 uint* d_output 一起工作，而到主机的副本被写入 uchar* h_output

uchar *h_output = (uchar*)malloc(width*height*depth);
checkCudaErrors(cudaMemcpy(h_output, d_output, width*height*depth*sizeof(uchar), cudaMemcpyDeviceToHost));

这导致了这种奇怪的行为