Function attributes / cache config / occupancy

Defined in code/simulator.h.

This page documents 7 symbols: 2 enums, 5 functions.

Enums

Signature	Description	Availability	Location	Example
enum cudaFuncAttribute		sim only	simulator.h:2336
enum cudaFuncCache		sim only	simulator.h:2334

Functions

Signature	Description	Availability	Location	Example
inline cudaError_t cudaFuncGetAttributes(cudaFuncAttributes* a, F )		sim only	simulator.h:2343	View
inline cudaError_t cudaFuncSetAttribute(F , cudaFuncAttribute , int )		sim only	simulator.h:2351	View
inline cudaError_t cudaFuncSetCacheConfig(F , cudaFuncCache )		sim only	simulator.h:2350	View
inline cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor( int* numBlocks, F , int , size_t )		sim only	simulator.h:2352	View
inline cudaError_t cudaOccupancyMaxPotentialBlockSize( int* minGridSize, int* blockSize, F , size_t = 0, int = 0)		sim only	simulator.h:2354

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

cudaFuncGetAttributes (test_web_cs_cppoverload.cu:138)

    void (*func1)(const int *, int *, int);
    void (*func2)(const int2 *, int *, int);
    void (*func3)(const int *, const int *, int *, int);
    struct cudaFuncAttributes attr;

    // overload 1: simple_kernel(const int*, int*, int)
    func1 = simple_kernel;
    memset(&attr, 0, sizeof(attr));
    CHECK(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared));
    CHECK(cudaFuncGetAttributes(&attr, *func1));
    OUTPUT_ATTR(attr);
    LAUNCH(func1, DIV_UP(N, THREAD_N), THREAD_N, dInput, dOutput, a);
    CHECK(cudaMemcpy(hOutput.data(), dOutput, sizeof(int) * N, cudaMemcpyDeviceToHost));
    funcResult = check_func1(hInput.data(), hOutput.data(), a);
    printf("simple_kernel(const int *pIn, int *pOut, int a) %s\n\n",
           funcResult ? "PASSED" : "FAILED");
    testResult &= funcResult;

    // overload 2: simple_kernel(const int2*, int*, int)
    func2 = simple_kernel;
    memset(&attr, 0, sizeof(attr));

cudaFuncSetAttribute (test_llmc_fused_residual_forward_full.cu:314)

void fused_residual_forward5(floatX* residual, floatX* normed, floatX* mean, floatX* rstd,
                             const floatX* inp1, const floatX* inp2,
                             const floatX* weight, const floatX* bias,
                             int N, int C, const int block_size) {
    int block_y = block_size / 32;
    const int grid_size = ceil_div(N, block_y);
    size_t smem = (2 + block_y) * C * sizeof(floatX);

    // in order to use more than 48 KiB of smem, need to call cudaFuncSetAttribute
    cudaCheck(cudaGetLastError());
    auto status = cudaFuncSetAttribute(fused_residual_forward_kernel5,
                                       cudaFuncAttributeMaxDynamicSharedMemorySize, (int)smem);
    cudaGetLastError();
    cudaCheck(status);   // [SIM] no kernel-4 fallback here; kernel 4 lives in the sibling test
    LAUNCH_DYN(fused_residual_forward_kernel5, grid_size, dim3(32, block_y), smem,
               residual, normed, mean, rstd, inp1, inp2, weight, bias, N, C);
    cudaCheck(cudaGetLastError());
}

void fused_residual_forward6(floatX* residual, floatX* normed, floatX* mean, floatX* rstd,
                             const floatX* inp1, const floatX* inp2,
                             const floatX* weight, const floatX* bias,

cudaFuncSetCacheConfig (test_web_cs_cppoverload.cu:137)

    int a = 1;

    void (*func1)(const int *, int *, int);
    void (*func2)(const int2 *, int *, int);
    void (*func3)(const int *, const int *, int *, int);
    struct cudaFuncAttributes attr;

    // overload 1: simple_kernel(const int*, int*, int)
    func1 = simple_kernel;
    memset(&attr, 0, sizeof(attr));
    CHECK(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared));
    CHECK(cudaFuncGetAttributes(&attr, *func1));
    OUTPUT_ATTR(attr);
    LAUNCH(func1, DIV_UP(N, THREAD_N), THREAD_N, dInput, dOutput, a);
    CHECK(cudaMemcpy(hOutput.data(), dOutput, sizeof(int) * N, cudaMemcpyDeviceToHost));
    funcResult = check_func1(hInput.data(), hOutput.data(), a);
    printf("simple_kernel(const int *pIn, int *pOut, int a) %s\n\n",
           funcResult ? "PASSED" : "FAILED");
    testResult &= funcResult;

    // overload 2: simple_kernel(const int2*, int*, int)
    func2 = simple_kernel;

cudaOccupancyMaxActiveBlocksPerMultiprocessor (test_web_olcf_hw9_task2.cu:119)

    cudaMemset(d_odata, 0, tsize);
    cudaMalloc(&d_idxs, test_dsize * sizeof(unsigned));

    cudaDeviceProp prop{};
    cudaError_t err = cudaGetDeviceProperties(&prop, 0);
    if (err != cudaSuccess) { std::printf("cuda error: %s\n", cudaGetErrorString(err)); return 1; }
    if (prop.cooperativeLaunch == 0) { std::printf("cooperative launch not supported\n"); return 0; }
    std::printf("Device 0: %s — number of SMs = %d\n", prop.name, prop.multiProcessorCount);

    int numBlkPerSM = 0;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlkPerSM, my_remove_if<mytype>, nTPB, 0);
    std::printf("number of blocks per SM = %d\n", numBlkPerSM);

    // [SIM] cap the cooperative grid (see the header note); the kernel is a
    // grid-stride loop so any grid size gives the same answer.
    int nBlocks = numBlkPerSM * prop.multiProcessorCount;
    if (nBlocks > 8) nBlocks = 8;
    dim3 grid((unsigned)nBlocks), block(nTPB);
    std::printf("cooperative grid: %d blocks of %d threads\n", nBlocks, nTPB);

    bool ok = true;