僅存儲迭代次數最少的值

Question

我有N次迭代和L個存儲單元。每次迭代都將寫入0到（L-1）個存儲單元中的任何一個。

我想根據迭代次數在全局內存中存儲數據。假設迭代K和K-1都寫入相同的位置。在全局內存中的最終結果應該是K.線程K-1不應該覆蓋線程K的結果。

我試過了下面的解決方案。我爲每個內存位置創建一個具有最大迭代次數的數組，並初始化爲-1。然後我檢查迭代次數是否大於當前值。如果是，我存儲該值並更新最大迭代次數。

這是我目前的代碼，但它沒有給出更大迭代的正確結果。

#include<stdio.h> 
#include"cuda.h" 
__global__ void fun(int *A,int *maxIndex,int *index1,int *lock) 
{ 
        int threadid=blockIdx.x*blockDim.x+threadIdx.x; 
        int iteration_no=threadid; 
        int index=index1[threadid]; 
        int exitFromLoop=1; 
        while(exitFromLoop==1) 
        { 
          int flag=atomicCAS(&lock[index],0,1); 
        if(flag==0) 


          { 
            if(maxIndex[index]<iteration_no) 
            {  
              A[index]=89; 
              maxIndex[index]=iteration_no; 
              __threadfence(); 
            } 
            else 
            { 
            } 
            //__threadfence(); 
            lock[index]=0; 
            exitFromLoop=0; 
          } 
          else 
          { 

          } 
        } 

}

int main() 
{ 
    int A[10]={10,20,30,40,50,60,70,80,90,100}; 
    int maxIndex[10]={-1,-1,-1,-1,-1,-1,-1,-1,-1,-1}; 
    int lock[10]={0,0,0,0,0,0,0,0,0,0}; 
    int index[8192]; 
    srand(0); 
    for(int ii=0;ii<8192;ii++) 
    { 
      index[ii]=rand()%10; 
    } 
    int *index1; 
    int *A1,*maxIndex1; 
    int *lock1; 
    cudaMalloc((void**)&lock1,sizeof(int)*10); 
    cudaMalloc((void**)&A1,sizeof(int)*10); 
    cudaMalloc((void**)&index1,sizeof(int)*8192); 
    cudaMalloc((void**)&maxIndex1,sizeof(int)*10); 
    cudaMemcpy(A1,&A,sizeof(int)*10,cudaMemcpyHostToDevice); 
    cudaMemcpy(lock1,&lock,sizeof(int)*10,cudaMemcpyHostToDevice); 
    cudaMemcpy(maxIndex1,&maxIndex,sizeof(int)*10,cudaMemcpyHostToDevice); 
    cudaMemcpy(index1,&index,sizeof(int)*8192,cudaMemcpyHostToDevice); 
    fun<<<16,512>>>(A1,maxIndex1,index1,lock1); 
    cudaMemcpy(&A,A1,sizeof(int)*10,cudaMemcpyDeviceToHost); 
    cudaMemcpy(&maxIndex,maxIndex1,sizeof(int)*10,cudaMemcpyDeviceToHost); 
    printf("\nindex \n"); 
    for(int i=0;i<8192;i++) 
    { 
      printf("%d\n",index[i]); 
    } 

    for(int i=0;i<10;i++) 
    { 
      printf(" %d max is %d\n",A[i],maxIndex[i]); 
    } 
}                         

Answer 1

我認爲這可能是你追求的。

對於A陣列中的每個元件，存在包含最後線程更新A陣列中的元件的iteration_no的maxIndex陣列中相應的元件。如果當前線程的版本高於iteration_no，則maxIndex將更新爲當前線程的iteration_no，並且該線程更新A中的元素。

如果當前線程的iteration_no較低，則A元素不會更新，maxIndex中的iteration_no不會更新。

#include<stdio.h> 
#include"cuda.h" 

__global__ void fun(int *A,int *maxIndex,int *index) 
{ 
    int iteration_no=blockIdx.x*blockDim.x+threadIdx.x; 
    int i=index[iteration_no]; 
    if (atomicMax(maxIndex + i, iteration_no) < iteration_no) { 
    A[i] = 89; 
    } 
} 

int main() 
{ 
    int A[10] = {10,20,30,40,50,60,70,80,90,100}; 
    int maxIndex[10]={-1}; 
    int index[8192]; 
    srand(0); 
    for(int ii=0;ii<8192;ii++) 
    { 
    index[ii]=rand()%10; 
    } 
    int *index_d; 
    int *A_d,*maxIndex_d; 
    cudaMalloc((void**)&A_d,sizeof(int)*10); 
    cudaMalloc((void**)&index_d,sizeof(int)*8192); 
    cudaMalloc((void**)&maxIndex_d,sizeof(int)*10); 
    cudaMemcpy(A_d,&A,sizeof(int)*10,cudaMemcpyHostToDevice); 
    cudaMemcpy(maxIndex_d,&maxIndex,sizeof(int)*10,cudaMemcpyHostToDevice); 
    cudaMemcpy(index_d,&index,sizeof(int)*8192,cudaMemcpyHostToDevice); 

    fun<<<16,512>>>(A_d,maxIndex_d,index_d); 

    cudaMemcpy(&A,A_d,sizeof(int)*10,cudaMemcpyDeviceToHost); 
    cudaMemcpy(&maxIndex,maxIndex_d,sizeof(int)*10,cudaMemcpyDeviceToHost); 
    printf("\nindex \n"); 

    for(int i=0;i<8192;i++) { 
    printf("%d\n",index[i]); 
    } 

    for(int i=0;i<10;i++) { 
    printf(" %d max is %d\n",A[i],maxIndex[i]); 
    } 
}