“cudaOccupancyMaxActiveBlocksPerMultiprocessor”返回的随机占用值

Question

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

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

我的 GPU： NVIDIA GeForce GTX 1070

问题：

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

示例代码：

源.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 代码之前，我强烈建议您：

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

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

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

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

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

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

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