cudaDeviceSynchronize에서 잘못된 메모리 액세스

Question

특정 크기의 Heat 2D 시뮬레이션을 실행할 때 '잘못된 메모리 액세스'오류가 발생한다는 매우 이상한 버그가 발생하지만 동일한 시뮬레이션을 실행하면 시뮬레이션이 잘 실행됩니다. 더 적은 요소로.

배열 크기를 늘리면이 예외가 발생하는 이유가 있습니까? Titan Black GPU (6GB 메모리)를 사용하고 있지만, 실행중인 시뮬레이션은 그 크기에 가깝지 않습니다. 4000 x 4000 시뮬레이션을 실행할 수 있다고 계산했지만 250 x 250을 초과하면 오류가 발생합니다.

장치에서 시뮬레이션 개체 배열을 인스턴스화 한 직후 오류가 발생합니다. 인스턴스화 코드는 다음과 같습니다.

template<typename PlaceType, typename StateType> 
__global__ void instantiatePlacesKernel(Place** places, StateType *state, 
     void *arg, int *dims, int nDims, int qty) { 
    unsigned idx = blockDim.x * blockIdx.x + threadIdx.x; 

    if (idx < qty) { 
     // set pointer to corresponding state object 
     places[idx] = new PlaceType(&(state[idx]), arg); 
     places[idx]->setIndex(idx); 
     places[idx]->setSize(dims, nDims); 
    } 
} 

template<typename PlaceType, typename StateType> 
Place** DeviceConfig::instantiatePlaces(int handle, void *argument, int argSize, 
     int dimensions, int size[], int qty) { 

    // add global constants to the GPU 
    memcpy(glob.globalDims,size, sizeof(int) * dimensions); 
    updateConstants(glob); 

    // create places tracking 
    PlaceArray p; // a struct to track qty, 
    p.qty = qty; 

    // create state array on device 
    StateType* d_state = NULL; 
    int Sbytes = sizeof(StateType); 
    CATCH(cudaMalloc((void**) &d_state, qty * Sbytes)); 
    p.devState = d_state; // save device pointer 

    // allocate device pointers 
    Place** tmpPlaces = NULL; 
    int ptrbytes = sizeof(Place*); 
    CATCH(cudaMalloc((void**) &tmpPlaces, qty * ptrbytes)); 
    p.devPtr = tmpPlaces; // save device pointer 

    // handle arg if necessary 
    void *d_arg = NULL; 
    if (NULL != argument) { 
     CATCH(cudaMalloc((void**) &d_arg, argSize)); 
     CATCH(cudaMemcpy(d_arg, argument, argSize, H2D)); 
    } 

    // load places dimensions 
    int *d_dims; 
    int dimBytes = sizeof(int) * dimensions; 
    CATCH(cudaMalloc((void**) &d_dims, dimBytes)); 
    CATCH(cudaMemcpy(d_dims, size, dimBytes, H2D)); 

    // launch instantiation kernel 
    int blockDim = (qty - 1)/BLOCK_SIZE + 1; 
    int threadDim = (qty - 1)/blockDim + 1; 
    Logger::debug("Launching instantiation kernel"); 
    instantiatePlacesKernel<PlaceType, StateType> <<<blockDim, threadDim>>>(tmpPlaces, d_state, 
      d_arg, d_dims, dimensions, qty); 
    CHECK(); 

    CATCH(cudaDeviceSynchronize()); // ERROR OCCURS HERE 

    // clean up memory 
    if (NULL != argument) { 
     CATCH(cudaFree(d_arg)); 
    } 
    CATCH(cudaFree(d_dims)); 
    CATCH(cudaMemGetInfo(&freeMem, &allMem)); 

    return p.devPtr; 
} 

이 코드는 충분히 작은 시뮬레이션에서 오류없이 실행되므로 작동하는 사용자 정의 유형을 가정하십시오. 크기가 250 x 250 요소를 초과하면 커널 함수의 위치와 상태 배열에있는 요소의 수가 오류가되는 것으로 나타났습니다. 어떤 통찰력도 대단 할 것입니다.

감사합니다.

Answer 1

너무 많은 메모리를 할당하고 있기 때문에 커널에 new이 실패했을 가능성이 높습니다.

커널 내부 new의 동작 및 제한은 in-kernel malloc입니다. 이러한 할당은 장치 힙으로 제한되며, 기본값은 8MB입니다. 250x250 배열 크기가 해당 범위 (8MB)에 해당하면 그 이상으로 이동하면 새 작업의 일부가 "자동으로"실패 (null 포인터를 반환)합니다. 그런 다음 해당 null 포인터를 사용하려고하면 불법적 인 메모리 액세스가 발생합니다.

몇 가지 권고 사항 : 당신이 new을 사용하는 커널에 문제가있는 경우 cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

를 사용

1. 당신이 필요로 얼마나 많은 공간이 그림 밖으로, 앞서 장치 힙에 시간의 사전 예약을 또는 malloc 일 경우, 디버그 목적으로는 아마도 디버그 매크로를 사용하여 NULL에 대해 반환 된 포인터를 검사하는 것이 유용 할 수 있습니다. 이것은 일반적으로 좋은 습관입니다.
2. here 설명 된 방법을 사용하여 더 명확하게 (특정 커널의 특정 줄로 로컬 화하는) 불법 메모리 액세스를 디버깅하는 방법을 배울 수 있습니다.