CUDAのthrustライブラリを使用したプログラムで "trivial_device_copy D->H failed" というエラーが出る

Question

前提・実現したいこと

char型配列A, int型配列Cに対して、以下の条件を満たすchar型配列Bを並列処理で作成したいです。

しかし、以下のようなエラーが出てしまい、エラーの原因がわからず困っています。
自分なりに調べてみましたが、自分と似たようなケースを見つけることができず、行き詰まっている状態です。
このエラーの原因と対処法を教えてください。
よろしくお願いします。

発生している問題・エラーメッセージ

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  trivial_device_copy D->H failed: cudaErrorIllegalAddress: an illegal memory access was encountered

該当のソースコード

#include <stdio.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void test(char *A, char *B, int *C,int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    B[i] = C[i] != 0 ? A[C[i] - 1] : '$';
}

int main () {
    const int n = 12;
    char A[] = "mississippi$";
    thrust::host_vector<char> h_B(n);
    thrust::device_vector<char> d_B = h_B;
    int C[] = {11,10,7,4,1,0,9,8,6,3,5,2};

    dim3 block(8, 1);
    dim3 grid((n + block.x - 1) / block.x, 1);
    char *pd_B = thrust::raw_pointer_cast(&d_B[0]);

    test<<< grid, block >>>(A, pd_B, C, n);

    h_B = d_B;  // エラー発生個所？

    for (int i = 0; i < n; i++) {
        printf("%c ", h_B[i]);
    }
}

Answer 1

#include <stdio.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void test(char* A, char* B, int* C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    B[i] = C[i] != 0 ? A[C[i] - 1] : '$';
}

int main() {
    const int n = 12;
    char A[] = "mississippi$";
    int C[] = { 11,10,7,4,1,0,9,8,6,3,5,2 };

    thrust::device_vector<char> d_A(A, A+n);
    thrust::device_vector<char> d_B(n);
    thrust::device_vector<int>  d_C(C, C+n);

    dim3 block(8, 1);
    dim3 grid((n + block.x - 1) / block.x, 1);
    test<<<grid, block >>>(
        thrust::raw_pointer_cast(&d_A[0]),
        thrust::raw_pointer_cast(&d_B[0]),
        thrust::raw_pointer_cast(&d_C[0]),
        n);

    thrust::host_vector<char> h_B = d_B;

    for (int i = 0; i < n; i++) {
        printf("%c ", h_B[i]);
    }
}

Comments

[episteme]

元コードのどこがマズかったか、わかりました?

[m3Zi8VHu]

回答していただき本当にありがとうございました。また、返信に気づかず返事が遅れてしまい、申し訳ありません。
元コードのエラーの原因はtest()関数に与える引数をすべてデバイス用にキャストしていなかったところであると予想していますが、正しいでしょうか？

[episteme]

キャストもなにも、A,C はデバイス・メモリ上にないのでtestの引数に渡したら失敗するぢゃろ、てことで。

[m3Zi8VHu]

理解できました。ありがとうございました！