Управление памятью/указателями памяти CUDA в задачах классов

Question

Я борюсь с проблемой управления памятью. Я продолжаю получать «Unspecified start failure» при копировании результатов на хост.

Мой код довольно прост - он генерирует два uints в каждом потоке и умножает их. У меня есть класс для обеспечения Random Number:

class CuRandCuRandomNumberProvider : 
{ 
public: 
    CuRandCuRandomNumberProvider(dim3 numBlocks, dim3 threadsPerBlock); 
    CuRandCuRandomNumberProvider(dim3 numBlocks, dim3 threadsPerBlock, unsigned int seed); 
    __device__ unsigned int GetRandomNumber(); 
    ~CuRandCuRandomNumberProvider(); 
protected: 
    curandState * states; 
    __device__ bool IsPrime(unsigned int number); 
}; 

CuRandCuRandomNumberProvider::CuRandCuRandomNumberProvider(dim3 numBlocks, dim3 threadsPerBlock) 
{ 
    int numberOfThreads = threadsPerBlock.x * threadsPerBlock.y * numBlocks.x * numBlocks.y; 
    std::cout << numberOfThreads << std::endl; 
    cudaMalloc (&this->states, numberOfThreads*sizeof(curandState)); 
    setup_kernel <<< numBlocks, threadsPerBlock >>> (this->states, time(NULL)); 
} 

__device__ unsigned int CuRandCuRandomNumberProvider::GetRandomNumber() 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 
    register float r = curand_uniform(&this->states[offset]); 
    return 0 + ((double)UINT_MAX) * r; 
} 

setup_kernel хранится в файле заголовка и выглядит следующим образом:

__global__ void setup_kernel (curandState * state, unsigned long seed) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 
    curand_init (seed, offset, 0, &state[offset]); 
} 

Мое главное ядро ​​очень проста и выглядит следующим образом:

__global__ void InitKernel(uint3 * ptr, CuRandCuRandomNumberProvider * provider) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 

    ptr[offset].x = provider->GetRandomNumber(); 
    ptr[offset].y = provider->GetRandomNumber(); 
    ptr[offset].z = ptr[offset].x * ptr[offset].y; 
} 

Выполнение в главном, где последние cudaMemcpy вызывает проблемы:

uint3 * pqnD; 

uint3 * pqnH = (uint3*)malloc(sizeof(uint3) * numberOfThreads); 
memset(pqnH,0,sizeof(uint3) * numberOfThreads); 

HANDLE_ERROR(cudaMalloc((void**)&pqnD, sizeof(uint3) * numberOfThreads)); 

CuRandCuRandomNumberProvider * provider = new CuRandCuRandomNumberProvider(numBlocks, threadsPerBlock); 

InitKernel<<<numBlocks, threadsPerBlock>>>(pqnD, provider); 

HANDLE_ERROR(cudaMemcpy(pqnH, pqnD, sizeof(uint3) * numberOfThreads, cudaMemcpyDeviceToHost)); // this line causes error 

HANDLE_ERROR(cudaFree(pqnD)); 

Если я сделать все explicily, как:

uint3 * pqnD; 

uint3 * pqnH = (uint3*)malloc(sizeof(uint3) * numberOfThreads); 

memset(pqnH,0,sizeof(uint3) * numberOfThreads); 

HANDLE_ERROR(cudaMalloc((void**)&pqnD, sizeof(uint3) * numberOfThreads)); 

curandState * states; 

cudaMalloc (&states, numberOfThreads*sizeof(curandState)); 

setup_kernel <<< numBlocks, threadsPerBlock >>> (states, time(NULL)); 

CuRandCuRandomNumberProvider * provider = new CuRandCuRandomNumberProvider(numBlocks, threadsPerBlock, states); 


InitKernel2<<<numBlocks, threadsPerBlock>>>(pqnD, states); 

HANDLE_ERROR(cudaMemcpy(pqnH, pqnD, sizeof(uint3) * numberOfThreads, cudaMemcpyDeviceToHost)); 

HANDLE_ERROR(cudaFree(pqnD)); 

Где setup_kernel точно так же и InitKernel2 выглядит следующим образом:

__global__ void InitKernel2(uint3 * ptr, curandState * states) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 

    ptr[offset].x = GetRandomNumber(states); 
    ptr[offset].y = GetRandomNumber(states); 
    ptr[offset].z =  ptr[offset].x *  ptr[offset].y; 
} 

и GetRandomNumber является:

__device__ unsigned int GetRandomNumber(curandState * states) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 
    register float r = curand_uniform(&states[offset]); 
    return 0 + ((double)UINT_MAX) * r; 

} 

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

Пожалуйста, помогите :)!

Answer 1

Это незаконно:

CuRandCuRandomNumberProvider * provider = new CuRandCuRandomNumberProvider(numBlocks, threadsPerBlock); 

InitKernel<<<numBlocks, threadsPerBlock>>>(pqnD, provider); 

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

ptr[offset].x = provider->GetRandomNumber(); 

(в конечном счете, приводит к :)

register float r = curand_uniform(&this->states[offset]); 

является незаконным.

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

CuRandCuRandomNumberProvider provider(numBlocks, threadsPerBlock); 

в InitKernel:

__global__ void InitKernel(uint3 * ptr, CuRandCuRandomNumberProvider provider) // change 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 

    ptr[offset].x = provider.GetRandomNumber(); // change 
    ptr[offset].y = provider.GetRandomNumber(); // change 
    ptr[offset].z = ptr[offset].x * ptr[offset].y; 
} 

в CuRandCuRandomNumberProvider :: GetRandomNumber():

__device__ unsigned int CuRandCuRandomNumberProvider::GetRandomNumber() 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int offset = x + y * blockDim.x * gridDim.x; 
    register float r = curand_uniform(&(states[offset])); // change 
    return 0 + ((double)UINT_MAX) * r; 
} 

(и я тоже удалил деструктор прототип , поскольку это мешало.)