$ cat t429.cu
#include <stdint.h>
#include <stdio.h>

__global__ void t(const uint64_t *in){

  int position = threadIdx.x;
        uint16_t a = (in[position] >> 48);
        uint16_t b = (in[position] >> 32) & 0xffff;
        uint16_t g = (in[position] >> 16) & 0xffff;
        uint16_t r = in[position] & 0xffff;
  float fa = a;
  float fb = b;
  float fg = g;
  float fr = r;
  printf("%f\n", fa);
  printf("%f\n", fb);
  printf("%f\n", fg);
  printf("%f\n", fr);
}

int main(){

  uint64_t *data;
  cudaMalloc(&data, sizeof(uint64_t));
  cudaMemset(((unsigned char *)data)+0, 1, 2);
  cudaMemset(((unsigned char *)data)+2, 2, 2);
  cudaMemset(((unsigned char *)data)+4, 3, 2);
  cudaMemset(((unsigned char *)data)+6, 4, 2);
  t<<<1,1>>>(data);
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_61 -o t429 t429.cu
$ cuda-memcheck ./t429
========= CUDA-MEMCHECK
1028.000000
771.000000
514.000000
257.000000
========= ERROR SUMMARY: 0 errors
$