CUDA提供的函数评估Occupancy

#include <iostream>                                                                                                                                 
#include <cuda_runtime.h>
#include <cuda_occupancy.h>

// Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d[idx] = a[idx] * b[idx];
}
// Host code
int main()
{
    int numBlocks; // Occupancy in terms of active blocks
    int blockSize = 32;
    // These variables are used to convert occupancy to warps
    int device;
    cudaDeviceProp prop;
    int activeWarps;
    int maxWarps;
    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocks,
    MyKernel,
    blockSize,
    0);
    activeWarps = numBlocks * blockSize / prop.warpSize;
            maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;
    std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;
return 0;
}

其实occupancy可以用nvpprof来获取的；倒是可以从这个例子看下occupancy的计算过程。

NV提供的例子选择合适的blocks, threads

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

// Host code
int 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 a
    cudaOccupancyMaxPotentialBlockSize(
    &minGridSize,
    &blockSize,
    (void*)MyKernel,
    0,arrayCount);
    // Round up according to array size
    gridSize = (arrayCount + blockSize - 1) / blockSize;
    MyKernel<<<gridSize, blockSize>>>(array, arrayCount);
    cudaDeviceSynchronize();
    // If interested, the occupancy can be calculated with
    // cudaOccupancyMaxActiveBlocksPerMultiprocessor
    return 0;
}

感觉这些API使用的应该不多正常情况下，自己可以测试下不同的blocks和grids的比例。