CUDA - 여러 번 커널 호출

Question

쓰려고하는 CUDA 프로그램에 문제가 있습니다. 나는 약 524k 부동 소수점 값 (1.0)의 배열을 가지고 있고 나는 모든 값을 추가하는 감소 기법을 사용하고있다. 한 번만 실행하려면 문제가 제대로 작동하지만 실제로 커널을 여러 번 실행하여 결국 10 억 개가 넘는 값을 합산 할 수 있습니다.

524k 청크에서이 작업을 수행하는 이유는 gpu에서 약 1 백만 개가 넘을 때 항상 제로가 돌아옵니다. 카드의 메모리를 초과해서는 안되지만 항상 그 시점에서 실패합니다.

어쨌든 커널을 한 번만 돌릴 때 모든 것이 정상적으로 작동합니다. 즉, 아무 루핑도 문제가되지 않습니다. 내가 루프를 실행하면 0으로 돌아옵니다. 내가 어떤 곳을 벗어나고 있다고 의심하지만, 나는 그것을 알아낼 수 없다. 그것은 나를 미치게합니다.

#include <stdio.h> 
#include <stdlib.h> 
#include "cutil.h" 

#define TILE_WIDTH  512 
#define WIDTH   524288 
//#define WIDTH   1048576 
#define MAX_WIDTH  524288 

#define BLOCKS   WIDTH/TILE_WIDTH 

__global__ void PartSum(float * V_d) 
{ 
    int tx = threadIdx.x; 
    int bx = blockIdx.x; 

    __shared__ float partialSum[TILE_WIDTH]; 

    for(int i = 0; i < WIDTH/TILE_WIDTH; ++i) 
    { 
     partialSum[tx] = V_d[bx * TILE_WIDTH + tx]; 
     __syncthreads(); 


     for(unsigned int stride = 1; stride < blockDim.x; stride *= 2) 
     { 
     __syncthreads(); 
     if(tx % (2 * stride) == 0) 
      partialSum[tx] += partialSum[tx + stride]; 
     } 
    } 

    if(tx % TILE_WIDTH == 0) 
     V_d[bx * TILE_WIDTH + tx] = partialSum[tx]; 
} 

int main(int argc, char * argv[]) 
{ 
    float * V_d; 
    float * V_h; 
    float * R_h; 
    float * Result; 
    float * ptr; 

    dim3 dimBlock(TILE_WIDTH,1,1); 
    dim3 dimGrid(BLOCKS,1,1); 

    // Allocate memory on Host 
    if((V_h = (float *)malloc(sizeof(float) * WIDTH)) == NULL) 
    { 
     printf("Error allocating memory on host\n"); 
     exit(-1); 
    } 

    if((R_h = (float *)malloc(sizeof(float) * MAX_WIDTH)) == NULL) 
    { 
     printf("Error allocating memory on host\n"); 
     exit(-1); 
    } 

    // If MAX_WIDTH is not a multiple of WIDTH, this won't work 
    if(WIDTH % MAX_WIDTH != 0) 
    { 
     printf("The width of the vector must be a multiple of the maximum width\n"); 
     exit(-3); 
    } 

    // Initialize memory on host with 1.0f 
    ptr = V_h; 
    for(long long i = 0; i < WIDTH; ++i) 
    { 
     *ptr = 1.0f; 
     ptr = &ptr[1]; 
    } 

    ptr = V_h; 

    // Allocate memory on device in global memory 
    cudaMalloc((void**) &V_d, MAX_WIDTH*(sizeof(float))); 
    float Pvalue = 0.0f; 
    for(int i = 0; i < WIDTH/MAX_WIDTH; ++i) 
    { 


    if((Result = (float *) malloc(sizeof(float) * WIDTH)) == NULL) 
    { 
     printf("Error allocating memory on host\n"); 
     exit(-4); 
    } 

    for(int j = 0; j < MAX_WIDTH; ++j) 
    { 
     Result[j] = *ptr; 
     ptr = &ptr[1]; 
    } 

     ptr = &V_h[i*MAX_WIDTH]; 
     // Copy portion of data to device 
     cudaMemcpy(V_d, Result, MAX_WIDTH*(sizeof(float)), cudaMemcpyHostToDevice); 

     // Execute Kernel 
     PartSum<<<dimGrid, dimBlock>>>(V_d); 

     // Copy data back down to host 
     cudaMemcpy(R_h, V_d, MAX_WIDTH*(sizeof(float)), cudaMemcpyDeviceToHost); 

     for(int i = 0; i < MAX_WIDTH; i += TILE_WIDTH) 
     { 
     Pvalue += R_h[i]; 
     } 
printf("Pvalue == %f\n", Pvalue); 

    free(Result); 


    } 

// printf("WIDTH == %d items\n", WIDTH); 
// printf("Value: %f\n", Pvalue); 

    cudaFree(V_d); 
    free(V_h); 
    free(R_h); 
    return(1); 
} 

좋아, 나는 내가 함께 할 수있는 문제를 좁혀 것 같아요 : 어떤 도움에 감사드립니다

,

감사합니다, 여기에

알은 코드입니다 장치의 V_d. 나는 배열의 경계를 넘어서고 있다고 생각한다. 실제로 필요한 메모리 양의 2 배를 할당하면 예상되는 결과로 프로그램이 완료됩니다. 문제는 문제의 원인을 파악할 수 없다는 것입니다.

알

Answer 1

나는 내가 여기 처음 버그를 발견 생각 :

if(tx % TILE_WIDTH == 0) 
     V_d[bx * TILE_WIDTH + tx] = partialSum[tx]; 

의 범위는 텍사스 0-511이며이 그래서 경우 조건이 없을 것 512 도달하지 못했다 참된. (tx % (TILE_WIDTH-1) == 0) 인 경우 으로 작성할 수 있습니다.

Answer 2

먼저이 모양과 도움을 주신 모든 분들께 감사드립니다.

두 번째로 마침내 내가 잘못하고있는 것을 알아 냈습니다. BLOCKS은 WIDTH/TILE_WIDTH이 아니라 MAX_WIDTH/TILE_WIDTH으로 정의되어야합니다. 내 멍청한 실수.

다시 한번 감사드립니다.