# HG changeset patch # User Shohei KOKUBO # Date 1402474937 -32400 # Node ID faaea4e1ce1c5ba31668b95371d1ece40f2b47b9 # Parent 6fced32f85fd8e3783e3d5c305397c4e838ff14d minor change diff -r 6fced32f85fd -r faaea4e1ce1c example/cuda_fft/fft.cu --- a/example/cuda_fft/fft.cu Wed Jun 11 11:24:58 2014 +0900 +++ b/example/cuda_fft/fft.cu Wed Jun 11 17:22:17 2014 +0900 @@ -18,8 +18,7 @@ j >>= (32-m); - dst[nid*n+j].x = src[nid*n+gid].x; - dst[nid*n+j].y = src[nid*n+gid].y; + dst[nid*n+j] = src[nid*n+gid]; } __global__ void @@ -41,14 +40,12 @@ float2 xa, xb, xbxx, xbyy, wab, wayx, wbyx, resa, resb; - xa.x = x[a].x; - xa.y = x[a].y; - xb.x = x[b].x; - xb.y = x[b].y; + xa = x[a]; + xb = x[b]; xbxx.x = xbxx.y = xb.x; xbyy.x = xbyy.y = xb.y; - wab.x = (float)((unsigned int)w[l].x); + wab.x = (float)((unsigned int)w[l].x ^ (unsigned int)0x0); wab.y = (float)((unsigned int)w[l].y ^ (unsigned int)flag); wayx.x = (float)((unsigned int)wab.y ^ (unsigned int)0x80000000); @@ -63,10 +60,8 @@ resb.x = xa.x - xbxx.x*wab.x + xbyy.x*wbyx.x; resb.y = xa.y - xbxx.y*wab.y + xbyy.y*wbyx.y; - x[a].x = resa.x; - x[a].y = resa.y; - x[b].x = resb.x; - x[b].y = resb.y; + x[a] = resa; + x[b] = resb; } __global__ void @@ -140,7 +135,6 @@ unsigned int iid = ygid * n + xgid; unsigned int oid = xgid * n + ygid; - dst[oid].x = src[iid].x; - dst[oid].y = src[iid].y; + dst[oid] = src[iid]; } } diff -r 6fced32f85fd -r faaea4e1ce1c example/cuda_fft/main.cc --- a/example/cuda_fft/main.cc Wed Jun 11 11:24:58 2014 +0900 +++ b/example/cuda_fft/main.cc Wed Jun 11 17:22:17 2014 +0900 @@ -123,7 +123,7 @@ CUcontext context; cuCtxCreate(&context, CU_CTX_SCHED_SPIN, device); - printf("%u\n", cuModuleLoad(&module, "fft.ptx")); + cuModuleLoad(&module, "fft.ptx"); char* pgm_file = init(args, argv);