¿Bajo qué circunstancias debe utilizar elvolatile ¿Palabra clave con la memoria compartida de un kernel CUDA? Entiendo quevolatile le dice al compilador que nunca almacene en caché ningún valor, pero mi pregunta es sobre el comportamiento con una matriz compartida:

__shared__ float products[THREADS_PER_ACTION];

// some computation
products[threadIdx.x] = localSum;

// wait for everyone to finish their computation
__syncthreads();

// then a (basic, ugly) reduction:
if (threadIdx.x == 0) {
    float globalSum = 0.0f;
    for (i = 0; i < THREADS_PER_ACTION; i++)
        globalSum += products[i];
}

Lo necesitoproducts ser volátil en este caso? Solo se puede acceder a cada entrada de la matriz mediante un único subproceso, excepto al final, donde todo se lee mediante la hebra 0. ¿Es posible que el compilador pueda almacenar en caché toda la matriz, por lo que necesito que estévolatile¿O solo almacenará en caché elementos?

¡Gracias!