为了记录，这是家庭作业，因此请尽可能少或尽可能多地记住这一点。我们使用常量内存来存储一个“掩码矩阵”，该矩阵将用于在更大的矩阵上执行卷积。当我在主机代码中时，我使用 cudaMemcpyToSymbol() 将掩码复制到常量内存中。
我的问题是，一旦将其复制并启动我的设备内核代码，设备如何知道在哪里访问常量内存掩码矩阵。是否有我需要在内核启动时传递的指针。教授给我们的大多数代码都不应更改(没有指向掩码的指针)，但总有可能他犯了一个错误(尽管这很可能是我对某些内容的理解)
常量内存声明是否应该包含在单独的 kernel.cu 文件中？
我正在最小化代码以显示与常量内存有关的事情。因此，请不要指出是否未初始化某些内容。有这方面的代码，但目前并不关心。
主.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;
}
最佳答案
在“经典”CUDA 编译中，您必须在同一翻译单元中定义所有代码和符号(纹理、常量内存、设备函数)以及访问它们的任何主机 API 调用(包括内核启动、绑定(bind)到纹理、复制到符号)。这意味着，实际上，在同一个文件中(或通过同一个文件中的多个包含语句)。这是因为“经典”CUDA 编译不包括设备代码链接器。
自 CUDA 5 发布以来，有可能使用单独的编译模式并将不同的设备代码对象链接到支持它的架构上的单个 fatbinary 有效负载中。在这种情况下，您需要使用 extern 关键字声明任何 __constant__ 变量并只定义一次符号。
如果您不能使用单独的编译，那么通常的解决方法是在与内核相同的 .cu 文件中定义 __constant__ 符号，并包含一个小的主机包装函数，该函数仅调用 cudaMemcpyToSymbol 来设置有问题的 __constant__ 符号。您可能会对内核调用和纹理操作执行相同的操作。