cudaMalloc и cudaMemcpy не работают с вызовом ядра

Question

У меня есть уже инициализированный массив, который я пытаюсь использовать в каждом потоке вызова ядра (каждый поток использует другую часть массива, поэтому нет зависимостей). Я создаю массив и сохраняю память на устройстве, используя cudaMalloc, и массив копируется с хоста на устройство с использованием cudaMemcpy.

Я передаю указатель, возвращаемый cudaMalloc, на вызов ядра, который будет использоваться каждым потоком.

int SIZE = 100; 

int* data = new int[SIZE]; 
int* d_data = 0; 

cutilSafeCall(cudaMalloc(&d_data, SIZE * sizeof(int))); 
for (int i = 0; i < SIZE; i++) 
    data[i] = i; 

cutilSafeCall(cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice)); 

Этот код был снят с here. Для вызова ядра.

kernel<<<blocks, threads>>> (results, d_data); 

я следить за результатами каждого потока с помощью Result-структуру. Следующий код работает без ошибок.

__global__ void mainKernel(Result res[], int* data){ 
    int x = data[0]; 
} 

Но когда я задаю это значение res:

__global__ void mainKernel(Result res[], int* data){ 
    int threadId = (blockIdx.x * blockDim.x) + threadIdx.x; 

    int x = data[0]; 

    res[threadId].x = x; 
} 

возникает ошибка:

cudaSafeCall() Ошибка выполнения API в файле, строка 355: незаконный доступ к памяти был встречен.

Та же ошибка появляется и при любой операции, включающей использование этого указателя

__global__ void mainKernel(Result res[], int* data){ 
    int threadId = (blockIdx.x * blockDim.x) + threadIdx.x; 

    int x = data[0]; 

    if (x > 10) 
     res[threadId].x = 5; 
} 

Там нет никаких проблем с определением res. Присвоение любого другого значения res[threadId].x не дает мне никакой ошибки.

Это выход работает Cuda-MemCheck:

========= Недопустимый __global__ чтения размера 4
========= в 0x00000150 в mainKernel (Результат *, int *)
========= потоком (86,0,0) в блоке (49,0,0)
========= Адрес 0x13024c0000 is вне границ
========= Сохраненный хоста трассировку до точки входа драйвера во время запуска ядра
========= Хост Оправа:/USR/Lib/x86_64-линукс-Гну /libcuda.so.1 (cuLaunchKerne L + 0x2cd) [0x150d6d]
========= Хост Рама: ./ из [0x2cc4b]
========= хост-кадр: ./ из [0x46c23]
= ======== Host Frame: ./ out [0x3e37]
========= Host Frame: ./ out [0x3ca1]
========= Host Frame : ./ out [0x3cd6]
========= Host Frame: ./ out [0x39e9]
========= Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]
========= Host Frame :./ Из [0x31b9]

EDIT:

Это пример полного кода:

#include <stdio.h> 
#include <stdlib.h> 
#include <math.h> 
#include <string.h> 
#include <iostream> 
#include <assert.h> 

typedef struct  
{ 
    int x,y,z; 
} Result; 

__global__ void mainKernel(Result pResults[], int* dataimage) 
{ 

    int threadId = (blockIdx.x * blockDim.x) + threadIdx.x; 

    int xVal = dataimage[0]; 
    if (xVal > 10) 
     pResults[threadId].x = 5; 

} 

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

    int NUM_THREADS = 5*5; 

    int SIZE = 100; 

    int* data = new int[SIZE]; 
    int* d_data = 0; 

    cutilSafeCall(cudaMalloc(&d_data, SIZE * sizeof(int))); 
    for (int i = 0; i < SIZE; i++) 
     data[i] = i; 

    cutilSafeCall(cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice)); 

    unsigned int GPU_ID = 1; // not actually :-) 
    // unsigned int GPU_ID = cutGetMaxGflopsDeviceId() ; 
    cudaSetDevice(GPU_ID); 

    Result * results_GPU = 0; 
    cutilSafeCall(cudaMalloc(&results_GPU, NUM_THREADS * sizeof(Result))); 

    Result * results_CPU = 0; 
    cutilSafeCall(cudaMallocHost(&results_CPU, NUM_THREADS * sizeof(Result))); 

    mainKernel<<<5,5>>> (results_GPU, d_data); 

    cudaThreadSynchronize(); 

    cutilSafeCall(cudaMemcpy(results_CPU, results_GPU, NUM_THREADS * sizeof(Result),cudaMemcpyDeviceToHost)); 

    cutilSafeCall(cudaFree(results_GPU)); 
    cutilSafeCall(cudaFreeHost(results_CPU)); 
    cudaThreadExit(); 

} //() 

Answer 1

Ваша проблема заключается в этой последовательности вызовов:

cutilSafeCall(cudaMalloc(&d_data, SIZE * sizeof(int))); 
    for (int i = 0; i < SIZE; i++) 
     data[i] = i; 

    cutilSafeCall(cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice)); 

    unsigned int GPU_ID = 1; 
    cudaSetDevice(GPU_ID); 

    Result * results_GPU = 0; 
    cutilSafeCall(cudaMalloc(&results_GPU, NUM_THREADS * sizeof(Result))); 

    Result * results_CPU = 0; 
    cutilSafeCall(cudaMallocHost(&results_CPU, NUM_THREADS * sizeof(Result))); 

    mainKernel<<<5,5>>> (results_GPU, d_data); 

Что эффективно происходит то, что вы выделяете d_data и запускаете ядро ​​разных графических процессоров, а d_data не действующий на GPU, вы запускаете ядро.

подробно, потому что вы звоните cudaMalloc для d_data перед тем cudaSetDevice, вы выделения d_data на устройства по умолчанию, а затем явно выделяя results_GPU и запустить ядро ​​на устройстве 1. Очевидно, устройство 1 и устройство по умолчанию не являются тот же GPU (перечисление устройств обычно начинается с 0 во время выполнения).

Если изменить код так:

unsigned int GPU_ID = 1; 
    cutilSafeCall(cudaSetDevice(GPU_ID)); 

    cutilSafeCall(cudaMalloc(&d_data, SIZE * sizeof(int))); 
    for (int i = 0; i < SIZE; i++) 
     data[i] = i; 

    cutilSafeCall(cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice)); 

    Result * results_GPU = 0; 
    cutilSafeCall(cudaMalloc(&results_GPU, NUM_THREADS * sizeof(Result))); 

    Result * results_CPU = 0; 
    cutilSafeCall(cudaMallocHost(&results_CPU, NUM_THREADS * sizeof(Result))); 

    mainKernel<<<5,5>>> (results_GPU, d_data); 

то выберите устройство не по умолчанию, прежде чем какие-либо средства выделяются, проблема должна исчезнуть. Причина этого не происходит с очень простым ядром:

__global__ void mainKernel(Result res[], int* data){ 
    int x = data[0]; 
} 

просто, что компилятор CUDA выполняет очень агрессивные оптимизации по умолчанию, и потому результат считывания data[0] на самом деле не используется, все чтение может быть оптимизировано, и у вас останется пустое ядро, которое ничего не делает. Только когда результат загрузки из памяти используется в записи в память, код не будет оптимизирован во время компиляции. Вы можете подтвердить это самостоятельно, разобрав код, испускаемый компилятором, если вам интересно.

Обратите внимание, что есть способы сделать эту работу над многопроцессорными системами, которые поддерживали ее, путем однорангового доступа, но это должно быть явно сконфигурировано в вашем коде для этого средства, которое будет использоваться.