CUDA如何在设备内核中访问由主机代码声明的常量内存？

Question

记录一下，这是一份作业，考虑到这一点，请适度提供帮助。我们使用常量内存来存储“掩模矩阵”，以便对较大的矩阵执行卷积。当我在主机代码中时，我正在使用cudaMemcpyToSymbol()将掩模复制到常量内存中。

我的问题是，一旦这个被复制了并启动设备内核代码，设备怎么知道在哪里访问常量内存掩模矩阵。是否需要在内核启动时传入指针？教授给我们的大部分代码都不应该改变（没有传递掩码的指针），但总会有错误的可能性（尽管最可能是我对某些事情的理解错误）。

常量内存声明是否应该包含在单独的kernel.cu文件中？

我正在将代码最小化，只显示与常量内存有关的内容。因此，请不要指出某些东西未初始化等问题。有相关的代码，但现在不是关注的重点。

main.cu:

#include <stdio.h>
#include "kernel.cu"

__constant__ float M_d[FILTER_SIZE * FILTER_SIZE];

int main(int argc, char* argv[])
{

     Matrix M_h, N_h, P_h; // M: filter, N: input image, P: output image

    /* Allocate host memory */
    M_h = allocateMatrix(FILTER_SIZE, FILTER_SIZE);
    N_h = allocateMatrix(imageHeight, imageWidth);
    P_h = allocateMatrix(imageHeight, imageWidth);

    /* Initialize filter and images */
    initMatrix(M_h);
    initMatrix(N_h);


    cudaError_t cudda_ret = cudaMemcpyToSymbol(M_d, M_h.elements, M_h.height * M_h.width * sizeof(float), 0, cudaMemcpyHostToDevice);
    //char* cudda_ret_pointer = cudaGetErrorString(cudda_ret);

    if( cudda_ret != cudaSuccess){
        printf("\n\ncudaMemcpyToSymbol failed\n\n");
        printf("%s, \n\n", cudaGetErrorString(cudda_ret));
    }


    // Launch kernel ----------------------------------------------------------
    printf("Launching kernel..."); fflush(stdout);

    //INSERT CODE HERE
    //block size is 16x16
    //              \\\\\\\\\\\\\**DONE**
    dim_grid = dim3(ceil(N_h.width / (float) BLOCK_SIZE), ceil(N_h.height / (float) BLOCK_SIZE));
    dim_block = dim3(BLOCK_SIZE, BLOCK_SIZE);



    //KERNEL Launch

    convolution<<<dim_grid, dim_block>>>(N_d, P_d);

    return 0;
}

kernel.cu: 这是我不知道如何访问常量内存的地方。

//__constant__ float M_c[FILTER_SIZE][FILTER_SIZE];

__global__ void convolution(Matrix N, Matrix P)
{
    /********************************************************************
    Determine input and output indexes of each thread
    Load a tile of the input image to shared memory
    Apply the filter on the input image tile
    Write the compute values to the output image at the correct indexes
    ********************************************************************/

    //INSERT KERNEL CODE HERE

    //__shared__ float N_shared[BLOCK_SIZE][BLOCK_SIZE];


    //int row = (blockIdx.y * blockDim.y) + threadIdx.y;
    //int col = (blockIdx.x * blockDim.x) + threadIdx.x;

}

- NDEthos

---

为什么你不能将指向常量内存的指针作为参数传递给卷积函数呢？ - rasen58

Answer 1

在“传统”的CUDA编译中，您必须在同一翻译单元中定义所有代码和符号（纹理、常量内存、设备函数），以及任何访问它们的主机API调用（包括内核启动、绑定到纹理、复制到符号）。这意味着在同一文件中（或通过同一文件中的多个包含语句）有效地完成。这是因为“传统”的CUDA编译不包括设备代码链接器。

自CUDA 5发布以来，有可能在支持的体系结构上使用分离编译模式并将不同的设备代码对象链接到单个fatbinary负载中。在这种情况下，您需要使用extern关键字声明任何__constant__变量，并且只需定义该符号一次。

如果无法使用分离编译，则通常的解决方法是在与您的内核相同的.cu文件中定义__constant__符号，并包含一个小型主机包装器函数，该函数仅调用cudaMemcpyToSymbol来设置相关的__constant__符号。您可能会对内核调用和纹理操作采取相同的做法。

- talonmies

Answer 2

下面是一个“最小尺寸”的示例，展示了使用__constant__符号的方法。您不需要将任何指针传递给__global__函数。

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>

__constant__ float test_const;

__global__ void test_kernel(float* d_test_array) {
    d_test_array[threadIdx.x] = test_const;
}

#include <conio.h>
int main(int argc, char **argv) {

    float test = 3.f;

    int N = 16;

    float* test_array = (float*)malloc(N*sizeof(float)); 

    float* d_test_array;
    cudaMalloc((void**)&d_test_array,N*sizeof(float));

    cudaMemcpyToSymbol(test_const, &test, sizeof(float));
    test_kernel<<<1,N>>>(d_test_array);

    cudaMemcpy(test_array,d_test_array,N*sizeof(float),cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) printf("%i %f\n",i,test_array[i]);

    getch();
    return 0;
}

- Vitality

---

我可以在一个文件中很好地完成它。但是当我有一个kernel.cu和一个main.cu，并且在main中声明时，我无法在kernel.cu中访问它。掩码矩阵是在main中创建的，那么我该如何在kernel中引用它？当我尝试时，它说它未定义。如果我在kernal中也放置__constant__ float M_d [FILTER_SIZE * FILTER_SIZE];，我会收到已经定义的错误。 - NDEthos