同じ__constant__メモリを使用するCUDAホストとデバイス

Question

私は定数メモリを使用するデバイス/ホスト機能を持っています。デバイス上では正常に動作しますが、ホスト上ではこのメモリは初期化されていないようです。

#include <iostream> 
#include <stdio.h> 


const __constant__ double vals[2] = { 0.0, 1000.0 }; 

__device__ __host__ double f(size_t i) 
{ 
    return vals[i]; 
} 

__global__ void kern() 
{ 
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]); 
} 

int main() { 
    std::cerr << f(0) << " " << f(1) << std::endl; 
    kern<<<1, 2>>>(); 
    cudaThreadSynchronize(); 
} 

このプリントは（CC 2.0以上が必要です）

0 0 
vals[0] = 0.000000 
vals[1] = 1000.000000 

問題とどのように私は、デバイスとホストメモリ定数の両方が同時に初期化され得ることができるとは何ですか？

Answer 1

私はMurphEngineerの答えに関する私のコメントにCygnusX1が誤解しているので、おそらく自分の答えを投稿するべきです。これはシグナスと同じ結果を持っていますが、それは実際のコードの顔に、より柔軟性のある

__constant__ double dc_vals[2] = { 0.0, 1000.0 }; 
     const double hc_vals[2] = { 0.0, 1000.0 }; 

__device__ __host__ double f(size_t i) 
{ 
#ifdef __CUDA_ARCH__ 
    return dc_vals[i]; 
#else 
    return hc_vals[i]; 
#endif 
} 

：それはのために、あなたが一定の配列での実行時に定義された値を持つことができます私はを意味することは、このましたたとえば、__constant__配列でcudaMemcpyToSymbol/cudsaMemcpyFromSymbolのようなCUDA API関数を使用することができます。

より現実的な例：

#include <iostream> 
#include <stdio.h> 

__constant__ double dc_vals[2]; 
     const double hc_vals[2]; 

__device__ __host__ double f(size_t i) 
{ 
#ifdef __CUDA_ARCH__ 
    return dc_vals[i]; 
#else 
    return hc_vals[i]; 
#endif 
} 

__global__ void kern() 
{ 
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]); 
} 

int main() { 
    hc_vals[0] = 0.0; 
    hc_vals[1] = 1000.0; 

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice); 

    std::cerr << f(0) << " " << f(1) << std::endl; 
    kern<<<1, 2>>>(); 
    cudaThreadSynchronize(); 
} 

Answer 2

__constant__修飾子を使用すると、そのメモリがデバイスに明示的に割り当てられます。新しいCUDA Unified Addressingのもの（cudaMalloc（）とその友人に割り当てられたメモリでのみ動作します）でも、ホストからそのメモリにアクセスする方法はありません。 constを使って変数を修飾すると、「これは（...）への定数ポインタです」と表示されます。

これを行う正しい方法は、実際にはホスト上に1つ、デバイス上に1つの2つの配列を持つことです。ホスト配列を初期化し、cudaMemcpyToSymbol（）を使用して、実行時にデバイス配列にデータをコピーします。これを行う方法の詳細については、このスレッドを参照してください。http://forums.nvidia.com/index.php?showtopic=69724

Answer 3

私はそれが動作しない理由 MurphEngineerがうまくを説明したと思います。すぐにこの問題を解決するには

、あなたはharrismの考えに従うことができ、このような何か：

#ifdef __CUDA_ARCH__ 
#define CONSTANT __constant__ 
#else 
#define CONSTANT 
#endif 

const CONSTANT double vals[2] = { 0.0, 1000.0 }; 

この方法でデバイスのコンパイルはデバイス__constant__コンパイルを作成する一方、ホストコンパイルは、通常のホストのconst配列を作成します。

このトリックでは、CUDA APIを使用してcudaMemcpyToSymbol()などの機能を持つデバイスアレイにアクセスすることが難しい場合があります。