用的块间循环展开时，《CUDA C编程权威指南》书上P300页写了会减少读/写事务
但我自己写了一段代码，似乎并不能减少读/写事务，我的配置是3050Ti-LapTop，CUDA11.7
编译选项是：nvcc -gencode=arch=compute_35,code=\"sm_35,compute_35\" -O0 ../ReadOffset.cu -o ReadOffset
代码在一楼

#include"cuda_runtime.h"
#include<iostream>
#define off 24
template<typename scalar_t>
__global__ void readOffsetUnroll4(scalar_t* A, scalar_t* B, scalar_t* C,
const int n, int offset)
{
unsigned int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
unsigned int k = i + offset;
if (k + 3 * blockDim.x < n)
{
C[i] = A[k] + B[k];
C[i + blockDim.x] = A[k + blockDim.x] + B[k + blockDim.x];
C[i + blockDim.x * 2] = A[k + blockDim.x * 2] + B[k + blockDim.x * 2];
C[i + blockDim.x * 3] = A[k + blockDim.x * 3] + B[k + blockDim.x * 3];
}
}
template<typename scalar_t>
__global__ void readOffset(scalar_t* A, scalar_t* B, scalar_t* C,
const int n, int offset)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int k = i + offset;
volatile scalar_t* temp = C;
if (i < n)
{
temp[i] = A[k] + B[k];
}
}
void launchUnroll(int offset)
{
int N = 1 << off;
size_t size = N * sizeof(int);
int* a = new int[N];
int* b = new int[N];
for (int i = 0; i < N; ++i)
{
a[i] = i;
b[i] = i;
}
int* a_cuda, * b_cuda, * c_cuda;
cudaMalloc((void**)&a_cuda, size);
cudaMalloc((void**)&b_cuda, size);
cudaMalloc((void**)&c_cuda, size);
cudaMemcpy(a_cuda, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(b_cuda, b, size, cudaMemcpyHostToDevice);
dim3 threadSize(1024);
dim3 blockSize(((N + threadSize.x - 1) / threadSize.x + 3) / 4);
readOffsetUnroll4<<<blockSize, threadSize>>>(a_cuda, b_cuda, c_cuda, N, offset);
int* c = new int[N];
cudaMemcpy(c, c_cuda, size, cudaMemcpyDeviceToHost);
cudaFree(a_cuda);
cudaFree(b_cuda);
cudaFree(c_cuda);
for (int i = 0; i < 10; ++i)
printf("%d, ", c[i]);
printf("\n");
}
void launch(int offset)
{
int N = 1 << off;
size_t size = N * sizeof(int);
int* a = new int[N];
int* b = new int[N];
for (int i = 0; i < N; ++i)
{
a[i] = i;
b[i] = i;
}
int* a_cuda, * b_cuda, * c_cuda;
cudaMalloc((void**)&a_cuda, size);
cudaMalloc((void**)&b_cuda, size);
cudaMalloc((void**)&c_cuda, size);
cudaMemcpy(a_cuda, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(b_cuda, b, size, cudaMemcpyHostToDevice);
dim3 threadSize(1024);
dim3 blockSize((N + threadSize.x - 1) / threadSize.x);
readOffset<<<blockSize, threadSize>>>(a_cuda, b_cuda, c_cuda, N, offset);
int* c = new int[N];
cudaMemcpy(c, c_cuda, size, cudaMemcpyDeviceToHost);
cudaFree(a_cuda);
cudaFree(b_cuda);
cudaFree(c_cuda);
for (int i = 0; i < 10; ++i)
printf("%d, ", c[i]);
printf("\n");
}
int main()
{
launch(0);
launch(2);
launch(4);
launch(8);
launchUnroll(0);
launchUnroll(2);
launchUnroll(4);
launchUnroll(8);
}

难道是因为编译器自动优化了，可我明明关掉了优化呀

这是非展开的baseline，和循环展开后的kernel对比，其中sectors一点都没变