L2 cache residency control ptx instruction

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I am trying to use the L2 cache residency control CUDA API to accelerate my kernel program.

size_t num_megabytes_persistent_data{3};

    cudaDeviceProp device_prop{};
    int current_device{0};
    CHECK_ERROR(cudaGetDevice(&current_device));
    CHECK_ERROR(cudaGetDeviceProperties(&device_prop, current_device));
    std::cout << "GPU: " << device_prop.name << std::endl;
    std::cout << "L2 Cache Size: " << device_prop.l2CacheSize / 1024 / 1024
              << " MB" << std::endl;
    std::cout << "Max Persistent L2 Cache Size: "
              << device_prop.persistingL2CacheMaxSize / 1024 / 1024 << " MB"
              << std::endl;

    size_t size = min(int(device_prop.l2CacheSize * 0.75), device_prop.persistingL2CacheMaxSize);
    size_t const num_megabytes_persistent_cache{3};

    /* Set aside max possible size of L2 cache for persisting accesses */
    CHECK_ERROR(
        cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize,
                           num_megabytes_persistent_cache * 1024 * 1024));
    std::cout << "Set persistent L2 cache size to "
              << num_megabytes_persistent_cache << " MB" << std::endl;
    std::cout << "cudaLimitPersistingL2CacheSize: " << cudaLimitPersistingL2CacheSize << " MB"
              << std::endl;

    // potentially thrashing stream
    cudaStream_t stream_persistent_cache;
    cudaStreamAttrValue stream_attribute_thrashing;
    // stream_attribute_thrashing.accessPolicyWindow.base_ptr =
    //     reinterpret_cast<void*>(d_input_symbols);
    // stream_attribute_thrashing.accessPolicyWindow.base_ptr =
    //     reinterpret_cast<void*>(d_active_state_list);
    stream_attribute_thrashing.accessPolicyWindow.base_ptr =
        reinterpret_cast<void*>(d_reordered_csr_row_ptrs);

    stream_attribute_thrashing.accessPolicyWindow.num_bytes =
        num_megabytes_persistent_data * 1024 * 1024;
    stream_attribute_thrashing.accessPolicyWindow.hitRatio = 1.0;
    stream_attribute_thrashing.accessPolicyWindow.hitProp =
        cudaAccessPropertyPersisting;
    stream_attribute_thrashing.accessPolicyWindow.missProp =
        cudaAccessPropertyNormal;

    CHECK_ERROR(cudaStreamSetAttribute(
        stream_persistent_cache, cudaStreamAttributeAccessPolicyWindow,
        &stream_attribute_thrashing));

    // no thrashing
    cudaStreamAttrValue stream_attribute_non_thrashing{
        stream_attribute_thrashing};
    stream_attribute_non_thrashing.accessPolicyWindow.hitRatio =
        std::min(static_cast<double>(num_megabytes_persistent_cache) /
                     num_megabytes_persistent_data,
                 1.0);
    CHECK_ERROR(cudaStreamSetAttribute(
        stream_persistent_cache, cudaStreamAttributeAccessPolicyWindow,
        &stream_attribute_non_thrashing));


kernel <<<>>>(d_input_symbols, d_active_state_list, d_reordered_csr_row_ptrs);

However, I do not see any improvement when I change this binded data ptr to different input parameter buffers (d_input_symbols, d_active_state_list, d_reordered_csr_row_ptrs);

Basically, I am doing a SpMSpV-type computation.

I am wondering if I am using the L2 resicency technique correctly and how can I confirm if I am correct.

Is there any related PTX instruction that I can check in my compiled PTX program?