标签: gpu-constant-memory

我在这里介绍一些代码

__constant__ int array[1024];

__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = array[tId];
}

__global__ void kernel2(int *d_dst, int *d_src) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = d_src[tId];
}

int main(int argc, char **argv) {
   int *d_array;
   int *d_src;
   cudaMalloc((void**)&d_array, sizeof(int) * 1024);
   cudaMalloc((void**)&d_src, sizeof(int) * 1024);

   int *test = new int[1024];
   memset(test, 0, sizeof(int) * 1024);

   for (int i = 0; i < 1024; i++) {
     test[i] = …

我试图了解手写内核的每个CUDA线程的资源使用情况.

我把我的kernel.cu文件编译成了一个kernel.o文件nvcc -arch=sm_20 -ptxas-options=-v

我得到了以下输出

ptxas info    : Compiling entry function 'searchkernel(octree, int*, double, int, double*, double*, double*)' for 'sm_20'
ptxas info    : Function properties for searchkernel(octree, int*, double, int, double*, double*, double*)
    72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14]

看看上面的输出,这是正确的

• 每个CUDA线程使用46个寄存器？
• 没有寄存器溢出到本地内存？

我在理解输出方面也遇到了一些问题.

• 我的内核调用了很多c++filt函数.IS 72字节是__device__和__global__函数堆栈帧的内存总和？

• __device__和之间有什么区别0 byte spill stores …

问题

我使用常量内存准备了一个示例CUDA代码.我可以在cuda 4.2中成功运行它,但是当我使用CUDA 5编译时,我得到 "无效的设备符号".我在这里附加了示例代码.

代码

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

struct CParameter
{
    int A;  
    float B;
    float C;
    float D;
};

__constant__ CParameter * CONSTANT_PARAMETER;   
#define PARAMETER "CONSTANT_PARAMETER"

bool ERROR_CHECK(cudaError_t Status)
{
    if(Status != cudaSuccess)
    {
        printf(cudaGetErrorString(Status));
        return false;
    }   
    return true;
}

// Kernel that executes on the CUDA device
__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N)
  {
      a[idx] = CONSTANT_PARAMETER->A * a[idx]; …

根据"CUDA C编程指南",只有在多处理器常量高速缓存被命中时,常量内存访问才会受益(见第5.3.2.4节)¹.否则,对于半翘曲,可能存在比合并的全局存储器读取更多的存储器请求.那么为什么常量内存大小限制为64 KB？

还有一个问题是为了不要问两次.据我所知,在Fermi架构中,纹理缓存与L2缓存相结合.纹理使用是否仍然有意义或全局内存读取是否以相同的方式缓存？

¹ 恒定存储器(第5.3.2.4节)

常量存储空间驻留在设备存储器中,并缓存在F.3.1和F.4.1节中提到的常量高速缓存中.

对于计算能力为1.x的设备,对于warp的常量内存请求首先被分成两个请求,每个半warp一个,独立发出.

然后,请求被分成多个单独的请求,因为初始请求中存在不同的内存地址,吞吐量减少的因子等于单​​独请求的数量.

然后,在高速缓存命中的情况下,或者在设备存储器的吞吐量下,以恒定高速缓存的吞吐量来服务所得到的请求.

我想了解一下如何分配常量内存(使用CUDA 4.2).我知道总可用常量内存是64KB.但是什么时候这个内存实际分配在设备上？此限制是否适用于每个内核,cuda上下文或整个应用程序？

假设.cu文件中有几个内核,每个内核使用少于64K的常量内存.但总的常量内存使用量超过64K.是否可以按顺序调用这些内核？如果使用不同的流同时调用它们会发生什么？

如果有一个大型CUDA动态库,其中包含大量内核,每个内核使用不同数量的常量内存,会发生什么？

如果有两个应用程序需要超过可用常量内存的一半,会发生什么？第一个应用程序运行正常,但第二个应用程序什么时候会失败？在应用程序启动时,在cudaMemcpyToSymbol()调用或内核执行时？

我自己无法弄明白,确保内核中使用的内存不变的最佳方法是什么.http://stackoverflow...r-pleasant-way也有类似的问题.我正在使用GTX580并仅编译2.0功能.我的内核看起来像

__global__ Foo(const int *src, float *result) {...}

我在主机上执行以下代码:

cudaMalloc(src, size);
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice);
Foo<<<...>>>(src, result);

另一种方法是添加

__constant__ src[size];

到.cu文件,从内核中删除src指针并执行

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice);
Foo<<<...>>>(result);

这两种方式是等效的还是第一种不保证使用常量内存而不是全局内存？大小动态变化,所以第二种方式在我的情况下不方便.

当我们使用xptxas检查寄存器使用情况时,我们会看到如下内容:

ptxas info : Used 63 registers, 244 bytes cmem[0], 51220 bytes cmem[2], 24 bytes cmem[14], 20 bytes cmem[16]

我想知道目前是否有任何文件清楚地解释了cmem [x].将常数存储器分成多个存储体,总共有多少存储体,除了0,2,14,16之外的其他存储体用于什么？

作为旁注,@ njuffa(特别感谢你)之前在nvidia的论坛上解释过什么是0,2,14,16银行:

使用的常量存储器在常量程序"变量"(存储区1)中分区,加上编译器生成的常量(存储区14).

cmem [0]:内核参数

cmem [2]:用户定义的常量对象

cmem [16]:编译器生成的常量(其中一些可能对应源代码中的文字常量)

我试图将我的模拟参数设置在恒定的内存中,但没有运气(CUDA.NET).cudaMemcpyToSymbol函数返回cudaErrorInvalidSymbol.cudaMemcpyToSymbol中的第一个参数是字符串......它是符号名吗？我不明白它是如何解决的.任何帮助赞赏.

//init, load .cubin   
float[] arr = new float[1];
    arr[0] = 0.0f;
    int size = Marshal.SizeOf(arr[0]) * arr.Length;
    IntPtr ptr = Marshal.AllocHGlobal(size);
    Marshal.Copy(arr, 0, ptr, arr.Length);
    var error = CUDARuntime.cudaMemcpyToSymbol("param", ptr, 4, 0, cudaMemcpyKind.cudaMemcpyHostToDevice);

我的.cu文件包含

__constant__ float param;

工作方案

     cuda.LoadModule(Path.Combine(Environment.CurrentDirectory, "name.cubin"));            
 simParams = cuda.GetModuleGlobal("params");
 float[] parameters = new float[N]{...}             
 cuda.CopyHostToDevice<float>(simParams, parameters);

我无法绕过对CUDA恒定内存的限制.

1. 为什么我们不能__constant__在运行时分配内存？为什么我需要在具有接近全局范围的固定大小变量中进行编译？

2. 什么是常量内存实际加载或卸载？我明白cudaMemcpytoSymbol 用于加载特定的数组,但每个内核是否都使用自己的常量内存分配？相关,是否有绑定和解除绑定的成本类似于绑定纹理的旧成本(也就是说,使用纹理为每个内核调用添加了成本)？

3. 常量内存驻留在芯片上的哪个位置？

我主要对与Pascal和Volta有关的答案感兴趣.

我读了一些引用本地内存的CUDA文档.(这主要是早期文档.)设备属性报告本地mem大小(每个线程)."本地"记忆是什么意思？什么是'本地'记忆？"本地"记忆在哪里？如何访问"本地"内存？这是__device__记忆,不是吗？

设备属性还报告:全局,共享和常量mem大小.这些陈述是否正确: 全局记忆是__device__记忆.它具有网格范围和网格(内核)的生命周期. 常量内存是__device__ __constant__内存.它具有网格范围和网格(内核)的生命周期. 共享内存是__device__ __shared__内存.它具有单个块范围和该块(线程)的生命周期.

我在想共享内存是SM内存.即只有那个SM才能直接访问的内存.资源相当有限.SM是不是一次分配了一堆块？这是否意味着SM可以交错执行不同的块(或不是)？即运行阻止*A*线程直到它们停止.然后运行block*B*线程直到它们停止.然后再次换回阻止*A*线程.或者SM是否为块*A*运行一组线程,直到它们停止.然后交换另一组块*A*线程.此交换继续,直到块*A*用尽.然后,只有这样才能在块*B*上开始工作.因为共享记忆,我问.如果单个SM从2个不同的块交换代码,那么SM如何快速交换/分出共享内存块？(我认为后来的senerio是真的,并且没有交换共享内存空间.块*A*运行直到完成,然后块*B*开始执行.注意:块*A*可能是不同的内核比块*B*.)

我有一些函数可以在常量设备内存中加载变量并启动内核函数.我注意到第一次将一个变量加载到常量内存中时需要0.6秒,但是常量内存上的下一个加载非常快(0.0008秒).无论哪个函数是main中的第一个函数,都会出现此行为.下面是一个示例代码:

        __constant__ double res1;

        __global__kernel1(...) {...}

        void function1() {
            double resHost = 255 / ((double) size);
            CUDA_CHECK_RETURN(cudaMemcpyToSymbol(res1, &resHost, sizeof(double)));


            //prepare and launch kernel
        }

        __constant__ double res2;

        __global__kernel2(...) {...}

        void function2() {
            double resHost = 255 / ((double) size);
            CUDA_CHECK_RETURN(cudaMemcpyToSymbol(res2, &resHost, sizeof(double)));


            //prepare and launch kernel
        }

        int main(){
            function1(); //takes 0.6 seconds for loading
            function2(); // takes 0.0008 seconds for loading
            function1(); //takes 0.0008 seconds for loading

            return 0;
        }

为什么会这样？我可以避免吗？

我有以下代码从主机变量复制到__constant__CUDA中的变量

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

    int exit_code;

    if (argc < 4) {
        std::cout << "Usage: \n " << argv[0] << " <input> <output> <nColors>" << std::endl;
        return 1;
    }

    Color *h_input;
    int h_rows, h_cols;

    timer1.Start();
    exit_code = readText2RGB(argv[1], &h_input, &h_rows, &h_cols);
    timer1.Stop();
    std::cout << "Reading: " << timer1.Elapsed() << std::endl;

    if (exit_code != SUCCESS){
        std::cout << "Error trying to read file." << std::endl;
        return FAILURE;
    }

    CpuTimer timer1;
    GpuTimer timer2;
    float timeStep2 = 0, timeStep3 = 0; …