"cudaOccupancyMaxActiveBlocksPerMultiprocessor" 返回的随机占用值

Question

我正在尝试了解“cudaOccupancyMaxActiveBlocksPerMultiprocessor”方法的用法和好处。

我使用的是 sample program present on NVIDIA developer forum 的略微修改版本。 基本上，我要求用户提供数组的大小。

我的显卡： NVIDIA GeForce GTX 1070

问题：

程序返回的占用值非常随机。很多时候，程序returns相同输入数组大小的占用值不同，程序有什么问题吗？
如屏幕截图所示，如果用户传递数组 size=512，则占用值为“13”，而如果我在程序中直接设置N=512，那么占用值为“47”。为什么?
为什么用户提供的数组大小=1024 占用值=0？

示例代码：

Source.cpp

#include "kernel_header.cuh"

#include <algorithm>
#include <iostream>

using namespace std;

int main(int argc, char* argv[])
{
    int N;
    int userSize = 0;

    //ask size to user
    cout << "\n\nType the size of 1D Array: " << endl;
    cin >> userSize;

    N = userSize>0? userSize : 1024; //<<<<<<<<<<<<<<<-------PROBLEM

    int* array = (int*)calloc(N, sizeof(int));
    for (int i = 0; i < N; i++)
    {
        array[i] = i + 1;
        //cout << "i = " << i << " is " << array[i]<<endl;
    }

    launchMyKernel(array, N);

    free(array);


    return 0;
}

kernel_header.cuh

#ifndef KERNELHEADER
#define KERNELHEADER

void launchMyKernel(int* array, int arrayCount);

#endif

kernel.cu

#include "stdio.h"
#include "cuda_runtime.h"

__global__ void MyKernel(int* array, int arrayCount)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arrayCount)
    {
        array[idx] *= array[idx];
    }
}

void launchMyKernel(int* array, int arrayCount)
{
    int blockSize;   // The launch configurator returned block size 
    int minGridSize; // The minimum grid size needed to achieve the 
                     // maximum occupancy for a full device launch 
    int gridSize;    // The actual grid size needed, based on input size 

    cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,MyKernel, 0, 0);

    // Round up according to array size 
    gridSize = (arrayCount + blockSize - 1) / blockSize;

    MyKernel << < gridSize, blockSize >> > (array, arrayCount);

    cudaDeviceSynchronize();

    // calculate theoretical occupancy
    int maxActiveBlocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&maxActiveBlocks,
        MyKernel, blockSize,
        0);

    int device;
    cudaDeviceProp props;
    cudaGetDevice(&device);
    cudaGetDeviceProperties(&props, device);

    float occupancy = (maxActiveBlocks * blockSize / props.warpSize) /
        (float)(props.maxThreadsPerMultiProcessor /
            props.warpSize);


    printf("\n\nMax. Active blocks found: %d\nOur Kernel block size decided: %d\nWarp Size: %d\nNumber of threads per SM: %d\n\n\n\n", maxActiveBlocks
        , blockSize,
        props.warpSize,
        props.maxThreadsPerMultiProcessor);

    printf("Launched blocks of size %d. Theoretical occupancy: %f\n",
        blockSize, occupancy);
}

Answer 1

在向其他人寻求无法按您预期的方式工作的 CUDA 代码之前，我强烈建议您：

使用proper CUDA error checking
运行您的代码带有消毒剂，例如 cuda-memcheck 或 compute-sanitizer

即使您不理解结果，报告的信息也会对那些试图帮助您的人有用。

在你的情况下，你正在对你的内核做一些非法的事情。具体来说，您已将主机指针传递给它（calloc 返回的是主机指针）。您几乎不能在 CUDA 中使用这样的指针（即用于 CUDA 设备代码），这是基本的 CUDA 编程原则。要了解构造此类代码的一种方法，以便您的内核实际上可以做一些有用的事情，请参阅 vectorAdd CUDA 示例代码。

当您的内核尝试使用此主机指针时，它会进行非法访问。至少在我的例子中，当我为数据大小输入 2048 并实施适当的 CUDA 错误检查时，我观察到内核和 所有后续的 CUDA activity returns 错误代码，包括您对 cudaOccupancyMaxActiveBlocksPerMultiprocessor 的调用。这意味着，该调用没有按照您的预期进行，并且它 returns 的数据是垃圾。

这至少是您获得垃圾计算值的原因之一。

当我解决该问题时（例如，通过将 calloc 替换为对 cudaMallocManaged 的适当设计的调用），然后您的代码为我报告占用计算为 1.0，输入数据大小为512、1024 和 2048。因此，我看不到任何可变性，充其量，如果您仍有疑问，我认为您需要重述它们（在新问题中）。

我并不是说如果你解决了这个问题，一切都会好起来的。但是这个问题掩盖了任何进行有用分析的能力。

"cudaOccupancyMaxActiveBlocksPerMultiprocessor" 返回的随机占用值

Random Occupancy values returned by the "cudaOccupancyMaxActiveBlocksPerMultiprocessor"

c++

cuda

nvidia