CUDA: Using single thread per block works but using multiple threads per block gives error

I'm a beginner to CUDA and was experimenting with basic programs. I have a 1D array with elements counting down from 100 <99,98,...2,1,0> . My function basically takes an element 'n' at index 'i' and allots the element at index 'n' to index 'i' of a new array. So applying this to the mentioned array should return <0,1,2,....,97,98,99>. And it works, only if I specify the threads per block as 1.

#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<cuda.h>
#include<cuda_runtime.h>


__global__ void shuffle(int *arr1,int *arr2,int n){

    int i = threadIdx.x + blockDim.x+blockIdx.x;

    while(i<n){
        arr2[i] = arr1[arr1[i]];
        i += blockDim.x;
    }
}


int main(){

    int m=10,n=10;

    int num = m*n;
    int size = num*sizeof(int);

    int *arr1,*arr2;

    arr1 = (int*)malloc(size);
    arr2 = (int*)malloc(size);

    for(int i=num-1;i>=0;i--){
        arr1[i] = i;
    }
    

    int *d_arr1,*d_arr2;

    cudaMalloc(&d_arr1,size);
    cudaMalloc(&d_arr2,size);

    cudaMemcpy(d_arr1,arr1,size,cudaMemcpyHostToDevice);

    shuffle<<<100,1>>>(d_arr1,d_arr2,num);

    cudaMemcpy(arr2,d_arr2,size,cudaMemcpyDeviceToHost);

    long error = 0;

    printf("Num: %d\n",num);

    //Prints value if value is right else prints correct value and actual value
    for(int i=0;i<num;i++){
        if(arr2[i] != i){
            error+=1;
            printf("%d %d\n",i,arr2[i]);
        }
        else{printf("%d\n",arr2[i]);}
    }

    printf("Error: %ld\n",error);
}

If I try calling the function as shuffle<<<25,4>>> i.e 25 blocks and 4 threads/block, I get the wrong values at indexes 1,4-1. Similarly shuffle<<<25,5>>> gives wrong values at indexes 1,5-1.

Solution

The problem is in your calculation of the array index for each thread.

You need to change the following line in the shuffle kernel:

int i = threadIdx.x + blockDim.x + blockIdx.x;

To:

/*-------------------------------V----------*/
int i = threadIdx.x + blockDim.x * blockIdx.x;

I.e. use multiplication instead of addition for calculating the array index.
This is because each block consists of blockDim.x threads, and so blockDim.x * blockIdx.x will give you the base offset for this block.