mirror of
https://github.com/epasveer/seer.git
synced 2026-08-29 16:40:42 +08:00
48 lines
1.1 KiB
Plaintext
48 lines
1.1 KiB
Plaintext
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
|
|
// Simple 8-bit bit reversal Compute test
|
|
|
|
#define N 256
|
|
|
|
__global__ void bitreverse (void *data) {
|
|
|
|
unsigned int *idata = (unsigned int*)data;
|
|
extern __shared__ int array[];
|
|
|
|
array[threadIdx.x] = idata[threadIdx.x];
|
|
|
|
array[threadIdx.x] = ((0xf0f0f0f0 & array[threadIdx.x]) >> 4) | ((0x0f0f0f0f & array[threadIdx.x]) << 4);
|
|
array[threadIdx.x] = ((0xcccccccc & array[threadIdx.x]) >> 2) | ((0x33333333 & array[threadIdx.x]) << 2);
|
|
array[threadIdx.x] = ((0xaaaaaaaa & array[threadIdx.x]) >> 1) | ((0x55555555 & array[threadIdx.x]) << 1);
|
|
|
|
idata[threadIdx.x] = array[threadIdx.x];
|
|
}
|
|
|
|
int main(void) {
|
|
|
|
void* d = NULL;
|
|
int i;
|
|
unsigned int idata[N], odata[N];
|
|
|
|
for (i = 0; i < N; i++) {
|
|
idata[i] = (unsigned int)i;
|
|
}
|
|
|
|
cudaMalloc((void**)&d, sizeof(int)*N);
|
|
cudaMemcpy(d, idata, sizeof(int)*N, cudaMemcpyHostToDevice);
|
|
|
|
bitreverse<<<1, N, N*sizeof(int)>>>(d);
|
|
|
|
cudaMemcpy(odata, d, sizeof(int)*N, cudaMemcpyDeviceToHost);
|
|
|
|
for (i = 0; i < N; i++) {
|
|
printf("%u -> %u\n", idata[i], odata[i]);
|
|
}
|
|
|
|
cudaFree((void*)d);
|
|
|
|
return 0;
|
|
}
|
|
|