Hi, I’ve recently started learning CUDA and inline PTX, and I’m trying to get a better grasp of how data is cached for the different cache levels.
I found a post from a year back that was asking a similar question regarding the granularity for how data is cached, but I have more questions based upon my tests playing around with the caches since it doesn’t add up to me.
I’m not sure if I’ve misunderstood the statement from the CUDA programming guide, but from my testing I feel like every memory operation is always served with 32 byte memory transaction, no matter if it is cached in only L2 or in both L1 and L2, since it never caches 128 bytes of data.
I’ve written simple kernels purely with inline PTX, targeting different memory units to see how they behave.
When I’m loading from memory using ld.global.cg.u64 and then perform another same load but at an address with 8-, 16- or 24-byte offset from the initial load, I get a L2 cache hit. If I target an with 32-, 64- or 96-byte from the inital load, I get a L2 cache miss (this follows the CUDA C Programming Guide “whereas memory accesses that are cached in L2 only are serviced with 32-byte memory transactions”).
When I’m loading from memory using ld.global.ca.u64 and then perform another same load but at an address with 8-, 16- or 24-byte offset from the initial load, I get a L1 cache hit. If I target an with 32-, 64- or 96-byte from the inital load, I get a L1 cache miss (this does not follow the CUDA C Programming Guide “Memory accesses that are cached in both L1 and L2 are serviced with 128-byte memory transactions”).
I’m executing the kernels with only only block and one thread to make it easier to analyze and minimize interference. I’m using NCU and these metrics to see L1 and L2 cache hits:
l1text_sectors_pipe_lsu_mem_global_op_ld_lookup_hit.sum
l1text_sectors_pipe_lsu_mem_global_op_ld_lookup_miss.sum
lts__t_sectors_srcunit_tex_op_read_lookup_hit.sum
lts__t_sectors_srcunit_tex_op_read_lookup_miss.sum
What have I missunderstood?
My GPU has the Lovelace microarchitecture and I’m using CUDA version 13.1.
If it’s any help, here’s the code for the test of checking ld.cg. The version with the cache all identifier is the same, just swap them out.
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"
#include "iostream"
#include <stdio.h>
cudaError_t global2global(long long* h_times);
__global__ void global_2_global(uint32_t *d_data, long long *d_dummy){
asm volatile (
"{\n\t"
// Declare virtual registers
".reg .u64 trd<7>;\n\t\n\t" // This dynamically creates 7 64-bit temp registers
"mov.u64 trd0, %0;\n\t"
"mov.u64 trd1, %1;\n\t\n\t"
"ld.global.cg.u64 trd2, [trd1];\n\t"
"ld.global.cg.u64 trd3, [trd1+8];\n\t"
"ld.global.cg.u64 trd4, [trd1+16];\n\t"
"ld.global.cg.u64 trd5, [trd1+24];\n\t"
"ld.global.cg.u64 trd6, [trd1+32];\n\t\n\t"
"add.u64 trd3, trd3, trd2;\n\t" // Dummy adds to prevent DCE
"add.u64 trd4, trd4, trd3;\n\t"
"add.u64 trd5, trd5, trd4;\n\t"
"add.u64 trd6, trd6, trd5;\n\t\n\t"
// Store accumalated dummy adds to d_dummy[0]
"st.global.u64 [trd0], trd6;\n\t"
"}"
: // No outputs mapped to C++ variables
: "l"(d_dummy), "l"(d_data) // %0 and %1
: "memory"
);
}
int main()
{
long long h_times[2] = {0};
cudaError_t cudaStatus = global2global(h_times);
cudaStatus = cudaDeviceReset();
return 0;
}
cudaError_t global2global(long long* h_times){
uint32_t* d_data;
long long* d_dummy;
cudaError_t cudaStatus;
cudaStatus = cudaMalloc(&d_data, 64);
cudaStatus = cudaMalloc(&d_dummy, sizeof(long long));
cudaStatus = cudaMemset(d_data, 42, 64); // Initialize some data to device memory
// Launch with a single thread
global_2_global << <1, 1 >> > (d_data, d_dummy);
cudaDeviceSynchronize();
cudaMemcpy(h_times, d_dummy, sizeof(long long), cudaMemcpyDeviceToHost);
std::cout << "--- Cache Test Results---" << std::endl;
std::cout << "Random Dummy Value : " << h_times[0] << std::endl;
return cudaStatus;
}