Почему появляется сообщение об ошибке при использовании cuda-memcheck, но при прямом запуске программы ошибки нет

Question

Мой код выглядит следующим образом:

template<typename scalar_t>
__global__ void compute_square(const int num, const scalar_t* data, scalar_t* res, float* time) {
    clock_t start = clock();
    const int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i{tid}; i < num; i+=stride) {
        res[i] = data[i] * data[i];
    }

    clock_t finish = clock();
    time[tid] = (float)(finish - start) / CLOCKS_PER_SEC;
}


template<typename scalar_t>
__global__ void compute_sum(const int num, const scalar_t* data, scalar_t* sum) {
    const int tid = threadIdx.x + blockIdx.x * blockDim.x;
    extern __shared__ scalar_t shared[]; // dynamic allocated
    shared[threadIdx.x] = 0; // assign 0 to the aligned memory
    __syncthreads();
    if (tid < num) {
        shared[threadIdx.x] = data[tid];
    }
    __syncthreads();

    for (int s=1; s < blockDim.x; s*=2) {
        int idx = 2 * s * threadIdx.x;
        if (idx < blockDim.x) {
            shared[idx] += shared[idx + s];
        }
        __syncthreads();
    }

    if (threadIdx.x == 1) { // must be one thread, or will be added many times
        atomicAdd(sum, shared[0]);
    }
}
void test_cuda_sync() {
    cout << "test sync" << endl;
    int len{1000};
    vector<float> data(len);
    vector<float> res(len); 
    vector<float> time(len);
    float sum{0};
    std::iota(data.begin(), data.end(), 0);

    // allocate memory
    float *dev_data{nullptr}, *dev_res{nullptr}, *dev_time{nullptr};
    float *dev_sum{nullptr};
    cudaMalloc((void**)&dev_data, sizeof(float) * len);
    cudaMalloc((void**)&dev_res, sizeof(float) * len);
    cudaMalloc((void**)&dev_time, sizeof(float) * len);
    cudaMalloc((void**)&dev_sum, sizeof(float));

    // copy data to device
    cudaMemcpy(dev_data, data.data(), sizeof(float) * len, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_sum, &sum, sizeof(float), cudaMemcpyHostToDevice);

    // execute kernel function
    dim3 block(512);
    dim3 grid(std::min(4096, (int)std::ceil(len / 512.))); 
    compute_square<float><<<grid, block, 0>>>(len, dev_data, dev_res, dev_time);
    compute_sum<float><<<grid, block, 4096>>>(len, dev_data, dev_sum);

    // copy results back to host
    cudaMemcpy(&res[0], dev_res, sizeof(int) * len, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time[0], dev_time, sizeof(float) * len, cudaMemcpyDeviceToHost);
    cudaMemcpy(&sum, dev_sum, sizeof(int), cudaMemcpyDeviceToHost);
    vector<int> tmp_cpu(len);

    // free allocated memory
    cudaFree(dev_data);
    cudaFree(dev_res);
    cudaFree(dev_time);
    cudaFree(dev_sum);

    // check results
    cout << "cuda reduce sum is: \n";
    cout << sum << endl;
    sum = 0;
    for (auto &el : data)
        sum += el;
    cout << "cpu sum is: \n";
    cout << sum << endl;

    for (int i{0}; i < 10; ++i) {
        cout << res[i] << ", ";
    } cout << endl;
    cout << (float)time[0] << endl;

}

Проблема в том, что: когда я компилирую его и запускаю скомпилированный исполняемый файл main с ./main, ошибки нет, и выходной результат правильно, однако, когда я запускаю с cuda-memcheck, появляется много сообщений об ошибках: cuda-memcheck ./main, сообщения об ошибках:

=========
========= Invalid __global__ write of size 4
=========     at 0x000001e0 in void compute_square<float>(int, float const *, float*, float*)
=========     by thread (488,0,0) in block (1,0,0)
=========     Address 0x7fd2a6202fa0 is out of bounds
=========     Device Frame:void compute_square<float>(int, float const *, float*, float*) (void compute_square<float>(int, float const *, float*, float*) : 0x1e0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2fe) [0x282a4e]
=========     Host Frame:./main [0x11e09]
=========     Host Frame:./main [0x11e97]
=========     Host Frame:./main [0x481e5]
=========     Host Frame:./main [0x3e10]
=========     Host Frame:./main (_Z14test_cuda_syncv + 0x211) [0x4691]
=========     Host Frame:./main (main + 0x12) [0x3b52]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf0) [0x20830]
=========     Host Frame:./main (_start + 0x29) [0x3c39]
=========

В чем проблема с моим кодом, пожалуйста, и почему не проблема, когда я запускаю программу, но при запуске с cuda-memcheck появляется много сообщений об ошибках?

Answer 1

Размер вашего массива 1000, верно? Проверьте индексы в выводе memcheck -

=========     by thread (488,0,0) in block (1,0,0)

и как вы вычисляете индекс в compute_square

    const int tid = threadIdx.x + blockIdx.x * blockDim.x;

488 + 1 * 512 равно 1000, что является первым индексом за пределами.