2016-06-24 4 views
0

私はこのコードを実行して推力のあるstable_sortとユーザー定義の演算子を使ってIPを比較します。 このコードは50000未満のIPの配列で動作しますが、大きな配列に対してはメモリエラーが発生します。 メモリの場所のエラー:大きな配列とユーザー定義の比較演算子を使用すると推力:: 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); 
} 

エラーは次のとおりです: のMicrosoft C++例外 ここで私が使用したコードはありますか? 私はこのソートを達成するために、別の手法を使用して、マージする部品の分割やソートなどを行うべきですか?

+0

[タグ:c]ではありません。ここ –

答えて

3

最近の問題点は、カーネル内のmallocnewは、割り当て可能なデバイスヒープのサイズが制限されていることです。この制限を上げることができます。 the documentationをお読みください。

いくつかの他の提案:

  1. あなたは(第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を示さない。

私はLinux上で作業しているので、goto文を削除しました(明らかにウィンドウに戻ってきました)。私は、gotoを使用するよりも、異なるエラー処理プロセスを考え出すことをお勧めします。それ以外の理由がない場合は、推力構造に問題が発生します。

また、カーネル内のnewまたはmallocは「ゆっくり」です。あなたはおそらく、適切なサイズの単一のcudaMallocコールで、必要な割り当てを前もって行うことによって、より大きなサイズのためにこれをスピードアップすることができます。残念ながら、ダブルポインタ配列dev_sorted_Ipsを使用すると、これは複雑になります。 1つのポインタ配列にフラット化し、必要なサイズをcudaMallocで1回割り当て、カーネルで必要な配列のインデックスを作成して動作させることをお勧めします。このコードをプロファイルすると、並べ替え操作ではなく、ご使用のprepare_ips_listカーネルで長い時間(例:size = 1000000)の実行時間の大部分が消費されることがわかります。したがって、パフォーマンス改善のための努力の焦点がそこから始まるはずです。

+0

はGPUのデバイス情報は、次のとおりです。 デバイス名:のGeForce GT 740M メモリクロックレート(キロヘルツ):900000 メモリバス幅(ビット):64 ピークメモリバンド幅(GB /秒):14.400000 私が使用していますあなたがcudaプロジェクトを作成するときのデフォルトの例は同じことを使用しているので、エラー処理のためのgoto文。 –

+0

Windowsでこれを実行している場合、 'prepare_ips_list'カーネルは大きな' n'に対して長い時間(数秒以上)を要することにも注意してください。それに対処するための何もしていなければ、このカーネルの実行時間はおそらくWDDMタイムアウトを引き起こします。これは、システムトレイでのドライバの再起動に関する通知と、カーネルコールの後に挿入された 'cudaDeviceSynchronize()'で返されたエラーによって明らかになります。 –

+0

それに対処する何かをすることによって、あなたはどういう意味ですか? ここで問題となるのは、2次元配列の必要性がユーザー定義演算子に関係していることです。 この配列を平坦化し、推力を使って並べ替える方法はありますか –

関連する問題