CUDAのthrustライブラリを使用したプログラムで "trivial_device

前提・実現したいこと

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

該当のソースコード

CUDA
1#include <stdio.h>
2#include <thrust/host_vector.h>
3#include <thrust/device_vector.h>
4#include <cuda.h>
5#include <cuda_runtime.h>
6
7__global__ void test(char *A, char *B, int *C,int n) {
8    int i = blockIdx.x * blockDim.x + threadIdx.x;
9    if (i >= n) return;
10    B[i] = C[i] != 0 ? A[C[i] - 1] : '$';
11}
12
13int main () {
14    const int n = 12;
15    char A[] = "mississippi$";
16    thrust::host_vector<char> h_B(n);
17    thrust::device_vector<char> d_B = h_B;
18    int C[] = {11,10,7,4,1,0,9,8,6,3,5,2};
19
20    dim3 block(8, 1);
21    dim3 grid((n + block.x - 1) / block.x, 1);
22    char *pd_B = thrust::raw_pointer_cast(&d_B[0]);
23
24    test<<< grid, block >>>(A, pd_B, C, n);
25
26    h_B = d_B;  // エラー発生個所？
27
28    for (int i = 0; i < n; i++) {
29        printf("%c ", h_B[i]);
30    }
31}

行動規範の内容に同意します

回答1件

ベストアンサー

C++
1#include <stdio.h>
2#include <thrust/host_vector.h>
3#include <thrust/device_vector.h>
4#include <cuda.h>
5#include <cuda_runtime.h>
6
7__global__ void test(char* A, char* B, int* C, int n) {
8    int i = blockIdx.x * blockDim.x + threadIdx.x;
9    if (i >= n) return;
10    B[i] = C[i] != 0 ? A[C[i] - 1] : '$';
11}
12
13int main() {
14    const int n = 12;
15    char A[] = "mississippi$";
16    int C[] = { 11,10,7,4,1,0,9,8,6,3,5,2 };
17
18    thrust::device_vector<char> d_A(A, A+n);
19    thrust::device_vector<char> d_B(n);
20    thrust::device_vector<int>  d_C(C, C+n);
21
22    dim3 block(8, 1);
23    dim3 grid((n + block.x - 1) / block.x, 1);
24    test<<<grid, block >>>(
25        thrust::raw_pointer_cast(&d_A[0]),
26        thrust::raw_pointer_cast(&d_B[0]),
27        thrust::raw_pointer_cast(&d_C[0]),
28        n);
29
30    thrust::host_vector<char> h_B = d_B;
31
32    for (int i = 0; i < n; i++) {
33        printf("%c ", h_B[i]);
34    }
35}
36