2016-06-24 5 views
0

나는이 코드를 실행하여 추력 stable_sort 및 사용자 정의 연산자를 사용하여 IP의 큰 배열을 정렬하여 IP를 비교합니다. 이 코드는 50000 개 미만의 IP 배열에서 작동하지만 큰 배열에 대해서는 메모리 오류가 발생합니다. 메모리 위치 오류 : 큰 배열 및 사용자 정의 비교 연산자를 사용할 때 thrust :: stable_sort

이 얼마나 큰 배열이 문제를 해결하기 위해 메모리 위치에서 추력 :: 시스템 :: SYSTEM_ERROR : 내가 가진

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <thrust/sort.h> 
#include <stdio.h> 
#include <time.h> 
#include <device_functions.h> 
template<typename T> 
struct vector_less 
{ 
    typedef T first_argument_type; 
    typedef T second_argument_type; 
    typedef bool result_type; 
    __host__ __device__ bool operator()(const T &lhs, const T &rhs) const { 
     if (lhs[0] == rhs[0]) 
     if (lhs[1] == rhs[1]) 
     if (lhs[2] == rhs[2]) 
      return lhs[3] < rhs[3]; 
     else 
      return lhs[2] < rhs[2]; 
     else 
      return lhs[1] < rhs[1]; 
     else 
      return lhs[0] < rhs[0]; 
    } 
}; 

__global__ void prepare_ips_list(unsigned char ** dev_sorted_Ips, unsigned char * ip_b1, unsigned char * ip_b2, unsigned char * ip_b3, unsigned char * ip_b4, unsigned int searchedIpsSize) 
{ 
    int thread = threadIdx.x + blockIdx.x * blockDim.x; 
    if (thread < searchedIpsSize) 
    { 
     dev_sorted_Ips[thread] = new unsigned char[4]; 
     dev_sorted_Ips[thread][0] = ip_b1[thread]; 
     dev_sorted_Ips[thread][1] = ip_b2[thread]; 
     dev_sorted_Ips[thread][2] = ip_b3[thread]; 
     dev_sorted_Ips[thread][3] = ip_b4[thread]; 
    } 

} 


int main() 
{ 
    const int size = 1000000; 

    unsigned char * ip_b1 = new unsigned char[size]; 
    unsigned char * ip_b2 = new unsigned char[size];; 
    unsigned char * ip_b3 = new unsigned char[size];; 
    unsigned char * ip_b4 = new unsigned char[size];; 

    unsigned char * dev_ip_b1; 
    unsigned char * dev_ip_b2; 
    unsigned char * dev_ip_b3; 
    unsigned char * dev_ip_b4; 

    unsigned char ** dev_sortedIps; 

    for (int i = 0; i < size; i++) 
    { 
     ip_b1[i] = rand() % 240; 
     ip_b2[i] = rand() % 240; 
     ip_b3[i] = rand() % 240; 
     ip_b4[i] = rand() % 240; 
    } 

    cudaError_t cudaStatus; 
    cudaStatus = cudaSetDevice(0); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); 
     goto Error; 
    } 

    cudaStatus = cudaMalloc((void**)&dev_ip_b1, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b1, ip_b1, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b2, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b2, ip_b2, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b3, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b3, ip_b3, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b4, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
     goto Error; 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b4, ip_b4, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
     goto Error; 
    } 

    cudaStatus = cudaMalloc((void**)&dev_sortedIps, size * sizeof(unsigned char *)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
     goto Error; 
    } 

    int resetThreads = size; 
    int resetBlocks = 1; 
    if (size > 1024) 
    { 
     resetThreads = 1024; 
     resetBlocks = size/1024; 
     if (size % 1024 > 0) 
      resetBlocks++; 
    } 

    prepare_ips_list << <resetBlocks, resetThreads >> >(dev_sortedIps, dev_ip_b1, dev_ip_b2, dev_ip_b3, dev_ip_b4, size); 



    thrust::device_ptr<unsigned char *> sorted_list_ptr1(dev_sortedIps); 
    thrust::stable_sort(sorted_list_ptr1, sorted_list_ptr1 + size, vector_less<unsigned char *>()); 

    cudaStatus = cudaGetLastError(); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "launch failed: %s\n", cudaGetErrorString(cudaStatus)); 
     goto Error; 
    } 

    // cudaDeviceSynchronize waits for the kernel to finish, and returns 
    // any errors encountered during the launch. 
    cudaStatus = cudaDeviceSynchronize(); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching !\n", cudaStatus); 
     goto Error; 
    } 

    return 0; 

Error: 
    cudaFree(dev_ip_b1); 
    cudaFree(dev_ip_b2); 
    cudaFree(dev_ip_b3); 
    cudaFree(dev_ip_b4); 
    cudaFree(dev_sortedIps); 
} 

오류는 다음과 같습니다 : 마이크로 소프트 C++ 예외 여기 내가 사용하는 코드는? 병합 후 부품을 나누고 정렬하는 것과 같은 정렬을 위해 다른 기술을 사용해야합니까?

+0

[tag : c]가 아닙니다.여기 –

답변

3

근위 문제는 커널에서 mallocnew은 할당 할 수있는 장치 힙의 크기가 제한된다는 것입니다. 이 한도를 높일 수 있습니다. the documentation을 읽어보십시오.

다른 몇 가지 제안 :

  1. 당신은 (첫 번째 추력 호출 전) 커널 후 확인 오류를하고 있지 않습니다. 커널에서 오류 검사를 수행하면 커널이 실패하고 있음을 알 수 있으며 추력은 단순히 오류를보고하는 것입니다. 혼동을 피하십시오. CUDA 코드에 문제가있는 경우 언제든지 rigorous, proper cuda error checking을 수행하십시오.

  2. 좋은 사례로서, 적어도 디버깅을 목적으로 NULL에 대해 new 또는 malloc에 의한 포인터 반환을 테스트하는 것은 좋지 않습니다. API가 할당 실패가 발생했음을 알리는 방법입니다.

아래 코드는 입력 크기에 대한 장치 힙을 조정하여 근위 문제에 대한 가능한 해결 방법을 보여줍니다. 당신이 명령 줄 매개 변수로 원하는 크기를 전달하여 다양한 크기를 테스트 할 수 있습니다

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <thrust/sort.h> 
#include <stdio.h> 
#include <time.h> 
#include <stdlib.h> 
#include <device_functions.h> 
#include <assert.h> 

template<typename T> 
struct vector_less 
{ 
    typedef T first_argument_type; 
    typedef T second_argument_type; 
    typedef bool result_type; 
    __host__ __device__ bool operator()(const T &lhs, const T &rhs) const { 
     if (lhs[0] == rhs[0]) 
     if (lhs[1] == rhs[1]) 
     if (lhs[2] == rhs[2]) 
      return lhs[3] < rhs[3]; 
     else 
      return lhs[2] < rhs[2]; 
     else 
      return lhs[1] < rhs[1]; 
     else 
      return lhs[0] < rhs[0]; 
    } 
}; 

__global__ void prepare_ips_list(unsigned char ** dev_sorted_Ips, unsigned char * ip_b1, unsigned char * ip_b2, unsigned char * ip_b3, unsigned char * ip_b4, unsigned int searchedIpsSize) 
{ 
    int thread = threadIdx.x + blockIdx.x * blockDim.x; 
    if (thread < searchedIpsSize) 
    { 
     dev_sorted_Ips[thread] = new unsigned char[4]; 
     if (dev_sorted_Ips[thread] == NULL) assert(0); 
     dev_sorted_Ips[thread][0] = ip_b1[thread]; 
     dev_sorted_Ips[thread][1] = ip_b2[thread]; 
     dev_sorted_Ips[thread][2] = ip_b3[thread]; 
     dev_sorted_Ips[thread][3] = ip_b4[thread]; 
    } 

} 


int main(int argc, char *argv[]) 
{ 

    int size = 50000; 
    if (argc > 1) size = atoi(argv[1]); 
    int chunks = size/50000 + 1; 
    cudaError_t cudaStatus; 
    cudaStatus = cudaDeviceSetLimit(cudaLimitMallocHeapSize, 8000000 * chunks); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "set device heap limit failed!"); 
    } 
    unsigned char * ip_b1 = new unsigned char[size]; 
    unsigned char * ip_b2 = new unsigned char[size];; 
    unsigned char * ip_b3 = new unsigned char[size];; 
    unsigned char * ip_b4 = new unsigned char[size];; 

    unsigned char * dev_ip_b1; 
    unsigned char * dev_ip_b2; 
    unsigned char * dev_ip_b3; 
    unsigned char * dev_ip_b4; 

    unsigned char ** dev_sortedIps; 

    for (int i = 0; i < size; i++) 
    { 
     ip_b1[i] = rand() % 240; 
     ip_b2[i] = rand() % 240; 
     ip_b3[i] = rand() % 240; 
     ip_b4[i] = rand() % 240; 
    } 

    cudaStatus = cudaSetDevice(0); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); 
    } 

    cudaStatus = cudaMalloc((void**)&dev_ip_b1, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b1, ip_b1, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b2, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b2, ip_b2, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b3, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b3, ip_b3, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
    } 
    cudaStatus = cudaMalloc((void**)&dev_ip_b4, size * sizeof(unsigned char)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
    } 
    cudaStatus = cudaMemcpy(dev_ip_b4, ip_b4, size * sizeof(unsigned char), cudaMemcpyHostToDevice); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMemcpy failed!"); 
    } 

    cudaStatus = cudaMalloc((void**)&dev_sortedIps, size * sizeof(unsigned char *)); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaMalloc failed!"); 
    } 

    int resetThreads = size; 
    int resetBlocks = 1; 
    if (size > 1024) 
    { 
     resetThreads = 1024; 
     resetBlocks = size/1024; 
     if (size % 1024 > 0) 
      resetBlocks++; 
    } 

    prepare_ips_list << <resetBlocks, resetThreads >> >(dev_sortedIps, dev_ip_b1, dev_ip_b2, dev_ip_b3, dev_ip_b4, size); 

    cudaStatus = cudaDeviceSynchronize(); 
    if (cudaStatus != cudaSuccess){ 
     printf(" kernel fail\n"); 
     exit(0);} 

    thrust::device_ptr<unsigned char *> sorted_list_ptr1(dev_sortedIps); 
    thrust::stable_sort(sorted_list_ptr1, sorted_list_ptr1 + size, vector_less<unsigned char *>()); 

    cudaStatus = cudaGetLastError(); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "launch failed: %s\n", cudaGetErrorString(cudaStatus)); 
    } 

    // cudaDeviceSynchronize waits for the kernel to finish, and returns 
    // any errors encountered during the launch. 
    cudaStatus = cudaDeviceSynchronize(); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching !\n", cudaStatus); 
    } 

    return 0; 

} 

참고 : 또한 다른 두 제안을 해결할 수있는 방법을 보여줍니다. 나는 1000000까지 시험했는데, 잘 작동하는 것 같았다. 결국 충분한 크기의 문제가 발생하면 GPU에서 메모리가 부족해질 것입니다. 당신은 가지고있는 GPU를 나타내지 않습니다.

나는 리눅스에서 일하고 있기 때문에 goto 문을 제거했습니다 (분명히 창으로 다시 전환했습니다). 나는 당신이 goto를 사용하는 것보다 다른 오류 처리 방법을 생각해내는 것이 좋습니다. 다른 이유가 없다면 추력 구조에 어려움이 생깁니다.

또한 커널 내 new 또는 malloc은 일종의 "느림"입니다. 적절한 크기의 단일 cudaMalloc 전화를 사용하여 필요한 할당을 사전에 수행함으로써 더 큰 규모의 전화를 걸면이 속도를 크게 높일 수 있습니다. 불행하게도 이중 포인터 배열 dev_sorted_Ips을 사용하면이 작업이 복잡해집니다. 대신 하나의 포인터 배열을 평평하게하고, cudaMalloc을 통해 필요한 크기를 한 번 할당하고 커널에서 필요한 배열 인덱싱을 수행하여 작동하게하는 것이 좋습니다. 이 코드의 프로필을 작성하면 정렬 작업이 아닌 prepare_ips_list 커널에서 대용량 실행 시간 (예 : size = 1000000)의 대부분을 소비한다는 것을 알 수 있습니다. 따라서 성능 향상을위한 노력의 초점이 시작되어야합니다.

+0

는 GPU 장치 정보입니다 : 장치 이름 : 지포스 GT 740M 메모리 클럭 속도 (KHZ) : 900,000 메모리 버스 폭 (비트) : 64 피크 메모리 대역폭 (GB/s의) : 나는를 사용하고 14.400000 오류 처리를위한 goto 문은 cuda 프로젝트를 만들 때 기본 예제가 동일한 것을 사용하기 때문에 발생합니다. –

+0

또한 윈도우에서 이것을 실행하는 경우,'prepare_ips_list' 커널은 더 큰'n'에 대해 오랜 시간 (몇 초 이상)을 필요로합니다. 이 문제를 해결하기 위해 아무 것도하지 않으면이 커널 실행 시간이 WDDM 시간 초과를 유발할 가능성이 높습니다. 이것은 시스템 트레이에서의 드라이버 재시작에 대한 통지와 커널 호출 후에 삽입 된'cudaDeviceSynchronize()'에서 리턴 된 에러에 의해 분명해질 것입니다. –

+0

문제를 해결하기 위해 뭔가하는 것이 무엇을 의미합니까? 여기에 문제가 있습니다. 2 차원 배열의 필요성은 사용자 정의 연산자와 관련됩니다. 이 배열을 평평하게하고 추력을 사용하여 정렬하는 방법이 있습니다 –