programming - nvidia cuda geforce

cudaMemcpyFromSymbol en una variable__device__

Question

Estoy tratando de aplicar una función kernel en una variable __device__ , que, de acuerdo con las especificaciones, reside "en la memoria global"

#include <stdio.h>
#include "sys_data.h"
#include "my_helper.cuh"
#include "helper_cuda.h"
#include <cuda_runtime.h>
double X[10] = {1,-2,3,-4,5,-6,7,-8,9,-10};
double Y[10] = {0};
__device__ double DEV_X[10];
int main(void) {
checkCudaErrors(cudaMemcpyToSymbol(DEV_X, X,10*sizeof(double)));
vector_projection<double><<<1,10>>>(DEV_X, 10);
getLastCudaError("oops");
checkCudaErrors(cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double)));
return 0;
}

La función kernel vector_projection se define en my_helper.cuh siguiente manera:

template<typename T> __global__ void vector_projection(T *dx, int n) {
int tid;
tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < n) {
if (dx[tid] < 0)
dx[tid] = (T) 0;
}
}

Como puede ver, utilizo cudaMemcpyToSymbol y cudaMemcpyFromSymbol para transferir datos hacia y desde el dispositivo. Sin embargo, recibo el siguiente error:

CUDA error at ../src/vectorAdd.cu:19 code=4(cudaErrorLaunchFailure)
"cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double))" 

Nota al pie: Por supuesto, puedo evitar el uso de variables __device__ e ir por algo como esto que funciona bien; Solo quiero ver cómo hacer lo mismo (si es posible) con __device__ variables __device__ .

Actualización: la salida de cuda-memcheck se puede encontrar en http://pastebin.com/AW9vmjFs . Los mensajes de error que recibo son los siguientes:

========= Invalid __global__ read of size 8
========= at 0x000000c8 in /home/ubuntu/Test0001/Debug/../src/my_helper.cuh:75:void vector_projection<double>(double*, int)
========= by thread (9,0,0) in block (0,0,0)
========= Address 0x000370e8 is out of bounds

Answer 1

La raíz del problema es que no puede tomar la dirección de una variable de dispositivo en el código de host común :

vector_projection<double><<<1,10>>>(DEV_X, 10);
^

Aunque parece que se compila correctamente, la dirección real pasada es basura.

Para tomar la dirección de una variable de dispositivo en el código de host, podemos usar cudaGetSymbolAddress

Aquí hay un ejemplo trabajado que compila y funciona correctamente para mí:

$ cat t577.cu
#include <stdio.h>
double X[10] = {1,-2,3,-4,5,-6,7,-8,9,-10};
double Y[10] = {0};
__device__ double DEV_X[10];
template<typename T> __global__ void vector_projection(T *dx, int n) {
int tid;
tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < n) {
if (dx[tid] < 0)
dx[tid] = (T) 0;
}
}
int main(void) {
cudaMemcpyToSymbol(DEV_X, X,10*sizeof(double));
double *my_dx;
cudaGetSymbolAddress((void **)&my_dx, DEV_X);
vector_projection<double><<<1,10>>>(my_dx, 10);
cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double));
for (int i = 0; i < 10; i++)
printf("%d: %f/n", i, Y[i]);
return 0;
}
$ nvcc -arch=sm_35 -o t577 t577.cu
$ cuda-memcheck ./t577
========= CUDA-MEMCHECK
0: 1.000000
1: 0.000000
2: 3.000000
3: 0.000000
4: 5.000000
5: 0.000000
6: 7.000000
7: 0.000000
8: 9.000000
9: 0.000000
========= ERROR SUMMARY: 0 errors
$

Esta no es la única forma de abordar esto. Es legal tomar la dirección de una variable de dispositivo en el código del dispositivo, para que pueda modificar su kernel con una línea como esta:

T *dx = DEV_X;

y renunciar a pasar la variable del dispositivo como un parámetro de kernel. Como se sugiere en los comentarios, también puede modificar su código para usar la memoria unificada .

Con respecto a la comprobación de errores, si se desvía de la correcta comprobación de error de cuda y no tiene cuidado con sus desviaciones, los resultados pueden ser confusos. La mayoría de las llamadas API de cuda pueden, además de los errores que surgen de su propio comportamiento, devolver un error que resulta de alguna actividad asíncrona CUDA previa (generalmente llamadas al kernel).