¿Alguien puede proporcionar un código de muestra que demuestre el uso del punto flotante de 16 bits en cuda?

Cuda 7.5 admite variables de coma flotante de 16 bits. ¿Alguien puede proporcionar un código de muestra que demuestre su uso?

Hay algunas cosas que debe tener en cuenta por adelantado:

  1. Consulte los intrínsecos de media precisión.
  2. Tenga en cuenta que la mayoría o todos estos intrínsecos solo son compatibles con el código del dispositivo. (Sin embargo, @njuffa ha creado un conjunto de funciones de conversión utilizables por el host aquí )
  3. Tenga en cuenta que los dispositivos de capacidad informática 5.2 y siguientes no son compatibles nativamente con la aritmética de media precisión. Esto significa que cualquier operación aritmética que se realice debe realizarse en algún tipo admitido, como float . Los dispositivos de capacidad informática 5.3 (Tegra TX1, actualmente) y dispositivos presumiblemente futuros, admitirán operaciones aritméticas de media precisión “nativas”, pero actualmente están expuestas a través de __hmul intrínsecos como __hmul . Un __hmul intrínseco __hmul estará definido en los dispositivos que no admiten operaciones nativas.
  4. Debe incluir cuda_fp16.h en cualquier archivo en el que pretenda utilizar estos tipos e intrínsecos en el código del dispositivo.

Con los puntos anteriores en mente, aquí hay un código simple que toma un conjunto de cantidades float , las convierte a half cantidades y las escala por un factor de escala:

 $ cat t924.cu #include  #include  #define DSIZE 4 #define SCF 0.5f #define nTPB 256 __global__ void half_scale_kernel(float *din, float *dout, int dsize){ int idx = threadIdx.x+blockDim.x*blockIdx.x; if (idx < dsize){ half scf = __float2half(SCF); half kin = __float2half(din[idx]); half kout; #if __CUDA_ARCH__ >= 530 kout = __hmul(kin, scf); #else kout = __float2half(__half2float(kin)*__half2float(scf)); #endif dout[idx] = __half2float(kout); } } int main(){ float *hin, *hout, *din, *dout; hin = (float *)malloc(DSIZE*sizeof(float)); hout = (float *)malloc(DSIZE*sizeof(float)); for (int i = 0; i < DSIZE; i++) hin[i] = i; cudaMalloc(&din, DSIZE*sizeof(float)); cudaMalloc(&dout, DSIZE*sizeof(float)); cudaMemcpy(din, hin, DSIZE*sizeof(float), cudaMemcpyHostToDevice); half_scale_kernel<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(din, dout, DSIZE); cudaMemcpy(hout, dout, DSIZE*sizeof(float), cudaMemcpyDeviceToHost); for (int i = 0; i < DSIZE; i++) printf("%f\n", hout[i]); return 0; } $ nvcc -o t924 t924.cu $ cuda-memcheck ./t924 ========= CUDA-MEMCHECK 0.000000 0.500000 1.000000 1.500000 ========= ERROR SUMMARY: 0 errors $ 

Si estudia el código anterior, notará que, excepto en el caso de dispositivos cc5.3 y superiores, la aritmética se realiza como una operación de float regular . Esto es consistente con la nota 3 anterior.

Los puntos para llevar son los siguientes:

  1. En dispositivos de cc5.2 y siguientes, el half tipo de datos puede ser útil, pero principalmente como una optimización de almacenamiento (y, tal vez, una optimización de ancho de banda de memoria, ya que una carga vectorial de 128 bits podría cargar 8 half cantidades a la vez ) Por ejemplo, si tiene una neural network grande, y ha determinado que los pesos pueden tolerar que se almacenen como cantidades de media precisión (duplicando así la densidad de almacenamiento, o aproximadamente el doble del tamaño de la neural network que puede representarse en el espacio de almacenamiento de una GPU), entonces podría almacenar los pesos de la neural network como de media precisión. Luego, cuando necesite realizar un pase hacia adelante (inferencia) o un pase hacia atrás (entrenamiento), puede cargar los pesos en la memoria, convertirlos sobre la marcha (usando los elementos intrínsecos) para float cantidades, realizar la operación necesaria ( quizás incluyendo ajustar el peso debido al entrenamiento), luego (si es necesario) almacene el peso nuevamente como una cantidad half .
  2. Para cc5.3 y dispositivos futuros, si el algoritmo lo tolera , es posible realizar una operación similar a la anterior, pero sin conversión a float (y tal vez a la half ), sino dejando todos los datos en la half representación, y haciendo la aritmética necesaria directamente (usando, por ejemplo, __hmul o __hadd intrinsics).

Aunque no lo he demostrado aquí, el half tipo de datos es "utilizable" en el código de host. Con eso, quiero decir que puedes asignar almacenamiento para artículos de ese tipo y realizar, por ejemplo, operaciones de cudaMemcpy en él. Pero el código de host no sabe nada sobre la half tipo de datos (por ejemplo, cómo hacer aritmética en él, o imprimirlo, o escribir conversiones) y los intrínsecos no se pueden utilizar en el código de host. Por lo tanto, puede asignar almacenamiento para una gran variedad de half tipo de datos si lo desea (quizás para almacenar un conjunto de pesos de redes neuronales), pero solo podría manipular directamente esos datos con facilidad desde el código del dispositivo, no desde el código de host.

Algunos más comentarios:

  1. La biblioteca CUBLAS implementa un multiplicador matriz-matriz diseñado para trabajar directamente en la half datos. La descripción anterior debería proporcionar una idea de lo que está sucediendo "debajo del capó" para diferentes tipos de dispositivos (es decir, capacidades de cómputo).

  2. Una pregunta relacionada sobre el uso de la half en el empuje está aquí .