Skip to content

Commit fdcee2a

Browse files
angerercangererc
committed
Code for the cuFFT callback blog post
Initial commit
1 parent 03742f8 commit fdcee2a

4 files changed

Lines changed: 499 additions & 0 deletions

File tree

‎posts/cufft-callbacks/Makefile‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
all: callbacks no_callbacks
2+
3+
callbacks:
4+
nvcc -ccbin g++ -dc -m64 -o cufft_callbacks.o -c cufft_callbacks.cu
5+
nvcc -ccbin g++ -m64 -o cufft_callbacks cufft_callbacks.o -lcufft_static -lculibos
6+
7+
no_callbacks:
8+
nvcc -ccbin g++ -dc -m64 -o cufft_no_callbacks.o -c cufft_no_callbacks.cu
9+
nvcc -ccbin g++ -m64 -o cufft_no_callbacks cufft_no_callbacks.o -lcufft -lculibos
10+
11+
test:
12+
nvcc -ccbin g++ -dc -m64 -o test.o -c test.cu
13+
nvcc -ccbin g++ -m64 -o test test.o -lcufft_static -lculibos
14+
15+
clean:
16+
rm *.o
17+
rm cufft_callbacks
18+
rm cufft_no_callbacks

‎posts/cufft-callbacks/common.h‎

Lines changed: 127 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,127 @@
1+
2+
#define INPUT_SIGNAL_SIZE 1024
3+
#define BATCH_SIZE 1000
4+
#define COMPLEX_SIGNAL_SIZE (INPUT_SIGNAL_SIZE/2 + 1)
5+
#define ITERATIONS 100
6+
7+
////////////////////////////////////////////////////////////////////////////////
8+
// CUDA error checking
9+
////////////////////////////////////////////////////////////////////////////////
10+
#define checkCudaErrors(val) __checkCudaErrors__ ( (val), #val, __FILE__, __LINE__ )
11+
12+
template <typename T>
13+
inline void __checkCudaErrors__(T code, const char *func, const char *file, int line)
14+
{
15+
if (code) {
16+
fprintf(stderr, "CUDA error at %s:%d code=%d \"%s\" \n",
17+
file, line, (unsigned int)code, func);
18+
cudaDeviceReset();
19+
exit(EXIT_FAILURE);
20+
}
21+
}
22+
23+
////////////////////////////////////////////////////////////////////////////////
24+
// Helper functions
25+
////////////////////////////////////////////////////////////////////////////////
26+
__device__ __host__ inline cufftComplex ComplexMul(cufftComplex a, cufftComplex b)
27+
{
28+
cufftComplex c;
29+
c.x = a.x * b.x - a.y * b.y;
30+
c.y = a.x * b.y + a.y * b.x;
31+
return c;
32+
}
33+
34+
void initInputs(char *dataIn, cufftComplex *filter) {
35+
srand(42);
36+
37+
// Initalize the memory for the signal
38+
for (size_t i = 0; i < INPUT_SIGNAL_SIZE * BATCH_SIZE; ++i)
39+
{
40+
if(i % INPUT_SIGNAL_SIZE == 0) srand(42);
41+
float val = rand() / (float)RAND_MAX;
42+
dataIn[i] = (char)(127 * val);
43+
}
44+
45+
// Initialize correction vector
46+
for(size_t i = 0; i < COMPLEX_SIGNAL_SIZE; i++) {
47+
srand(42);
48+
filter[i].x = rand() / (float)RAND_MAX;
49+
filter[i].y = 0.5f;
50+
}
51+
}
52+
53+
////////////////////////////////////////////////////////////////////////////////
54+
// Verification
55+
////////////////////////////////////////////////////////////////////////////////
56+
bool postprocess(const cufftComplex *ref, const cufftComplex *res, int size)
57+
{
58+
bool passed = true;
59+
for (int i = 0; i < size; i++)
60+
if (res[i].x != ref[i].x || res[i].y != ref[i].y) {
61+
printf("%d: (%4.2f,%4.2f) != (%4.2f, %4.2f)\n", i, res[i].x, res[i].y, ref[i].x, ref[i].y);
62+
printf("%25s\n", "*** FAILED ***");
63+
passed = false;
64+
break;
65+
}
66+
return passed;
67+
}
68+
69+
//CPU Versions of the custom kernels
70+
void ConvertInputR_onCPU(
71+
const char * __restrict__ dataIn,
72+
cufftReal * __restrict__ dataOut,
73+
size_t size)
74+
{
75+
for(size_t i = 0; i < size; i++) {
76+
char element = ((char*)dataIn)[i];
77+
dataOut[i] = (cufftReal)((float)element/127.0f);
78+
}
79+
}
80+
81+
void ConvolveAndStoreTransposedC_onCPU(
82+
const cufftComplex * __restrict__ dataIn,
83+
cufftComplex * __restrict__ dataOut,
84+
const cufftComplex * __restrict__ filter)
85+
{
86+
for(size_t row = 0; row < BATCH_SIZE; row++) {
87+
for(size_t col = 0; col < COMPLEX_SIGNAL_SIZE; col++) {
88+
cufftComplex value = ComplexMul(dataIn[row * COMPLEX_SIGNAL_SIZE + col], filter[col]);
89+
dataOut[col * BATCH_SIZE + row] = value;
90+
}
91+
}
92+
}
93+
94+
/*
95+
* computes the reference and returns a newly allcoated cuda-managed cufftComplex* array.
96+
* The caller is responsible for freeing the reference array
97+
*/
98+
cufftComplex* computeReference(const char *dataIn, const cufftComplex *filter) {
99+
cufftReal *tmp_result1;
100+
cufftComplex *tmp_result2, *result;
101+
102+
checkCudaErrors(cudaMallocManaged(&tmp_result1, sizeof(cufftReal) * INPUT_SIGNAL_SIZE * BATCH_SIZE, cudaMemAttachGlobal));
103+
checkCudaErrors(cudaMallocManaged(&tmp_result2, sizeof(cufftComplex) * COMPLEX_SIGNAL_SIZE * BATCH_SIZE, cudaMemAttachGlobal));
104+
checkCudaErrors(cudaMallocManaged(&result, sizeof(cufftComplex) * COMPLEX_SIGNAL_SIZE * BATCH_SIZE, cudaMemAttachGlobal));
105+
106+
ConvertInputR_onCPU(dataIn, tmp_result1, INPUT_SIGNAL_SIZE * BATCH_SIZE);
107+
108+
//We use cuFFT to compute the reference; we want to verify the callbacks and custom kernels
109+
//and not the FFT itself, so that's fine
110+
cufftHandle fftPlan;
111+
size_t workSize;
112+
113+
checkCudaErrors(cufftCreate(&fftPlan));
114+
int signalSize = INPUT_SIGNAL_SIZE;
115+
checkCudaErrors(cufftMakePlanMany(fftPlan, 1, &signalSize, 0,0,0,0,0,0, CUFFT_R2C, BATCH_SIZE, &workSize));
116+
117+
checkCudaErrors(cufftExecR2C(fftPlan, tmp_result1, tmp_result2));
118+
checkCudaErrors(cudaDeviceSynchronize());
119+
120+
ConvolveAndStoreTransposedC_onCPU(tmp_result2, result, filter);
121+
122+
checkCudaErrors(cufftDestroy(fftPlan));
123+
checkCudaErrors(cudaFree(tmp_result1));
124+
checkCudaErrors(cudaFree(tmp_result2));
125+
126+
return result;
127+
}

‎posts/cufft-callbacks/cufft_callbacks.cu‎

Lines changed: 156 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,156 @@
1+
/*
2+
* Copyright 1993-2014 NVIDIA Corporation. All rights reserved.
3+
*
4+
* Please refer to the NVIDIA end user license agreement (EULA) associated
5+
* with this source code for terms and conditions that govern your use of
6+
* this software. Any use, reproduction, disclosure, or distribution of
7+
* this software and related documentation outside the terms of the EULA
8+
* is strictly prohibited.
9+
*
10+
*/
11+
12+
#include <stdlib.h>
13+
#include <stdio.h>
14+
#include <string.h>
15+
#include <math.h>
16+
#include <cuda_runtime.h>
17+
#include <cufft.h>
18+
#include <cufftXt.h>
19+
20+
#include "common.h"
21+
22+
#define TILE_DIM 32
23+
#define BLOCK_ROWS 8
24+
25+
////////////////////////////////////////////////////////////////////////////////
26+
// Callback Implementations
27+
////////////////////////////////////////////////////////////////////////////////
28+
__device__ cufftReal CB_ConvertInputR(void *dataIn, size_t offset, void *callerInfo, void *sharedPtr) {
29+
char element = ((char*)dataIn)[offset];
30+
return (cufftReal)((float)element/127.0f);
31+
}
32+
33+
__device__ cufftCallbackLoadR d_loadCallbackPtr = CB_ConvertInputR;
34+
35+
__device__ void CB_ConvolveAndStoreTransposedC(void *dataOut, size_t offset, cufftComplex element, void *callerInfo, void *sharedPtr) {
36+
cufftComplex *filter = (cufftComplex*)callerInfo;
37+
size_t row = offset / COMPLEX_SIGNAL_SIZE;
38+
size_t col = offset % COMPLEX_SIGNAL_SIZE;
39+
40+
((cufftComplex*)dataOut)[col * BATCH_SIZE + row] = ComplexMul(element, filter[col]);
41+
}
42+
43+
__device__ cufftCallbackStoreC d_storeCallbackPtr = CB_ConvolveAndStoreTransposedC;
44+
45+
////////////////////////////////////////////////////////////////////////////////
46+
// Program main
47+
////////////////////////////////////////////////////////////////////////////////
48+
int main(int argc, const char **argv)
49+
{
50+
struct cudaDeviceProp properties;
51+
int device = argc > 1 ? atoi(argv[1]) : 0;
52+
53+
checkCudaErrors(cudaGetDevice(&device));
54+
checkCudaErrors(cudaGetDeviceProperties(&properties, device));
55+
if( !(properties.major >= 2) ) {
56+
printf("This sample requires CUDA architecture SM2.0 or higher\n");
57+
exit(EXIT_FAILURE);
58+
}
59+
60+
// Allocate and initialize memory
61+
printf("Preparing input: %dx%d\n", BATCH_SIZE, INPUT_SIGNAL_SIZE);
62+
char *_8bit_signal;
63+
cufftComplex *result, *filter;
64+
65+
checkCudaErrors(cudaMallocManaged(&_8bit_signal, sizeof(char) * INPUT_SIGNAL_SIZE * BATCH_SIZE, cudaMemAttachGlobal));
66+
checkCudaErrors(cudaMallocManaged(&result, sizeof(cufftComplex) * COMPLEX_SIGNAL_SIZE * BATCH_SIZE, cudaMemAttachGlobal));
67+
checkCudaErrors(cudaMallocManaged(&filter, sizeof(cufftComplex) * COMPLEX_SIGNAL_SIZE, cudaMemAttachGlobal));
68+
69+
initInputs(_8bit_signal, filter);
70+
71+
//compute reference result for later verification
72+
printf("Computing reference solution\n");
73+
cufftComplex *reference = computeReference(_8bit_signal, filter);
74+
75+
printf("Creating FFT plan\n");
76+
cufftHandle fftPlan;
77+
size_t workSize;
78+
79+
checkCudaErrors(cufftCreate(&fftPlan));
80+
int signalSize = INPUT_SIGNAL_SIZE;
81+
checkCudaErrors(cufftMakePlanMany(fftPlan, 1, &signalSize, 0,0,0,0,0,0, CUFFT_R2C, BATCH_SIZE, &workSize));
82+
83+
/*
84+
* Retrieve address of callback functions on the device
85+
*/
86+
cufftCallbackLoadR h_loadCallbackPtr;
87+
cufftCallbackStoreC h_storeCallbackPtr;
88+
checkCudaErrors(cudaMemcpyFromSymbol(&h_loadCallbackPtr,
89+
d_loadCallbackPtr,
90+
sizeof(h_loadCallbackPtr)));
91+
checkCudaErrors(cudaMemcpyFromSymbol(&h_storeCallbackPtr,
92+
d_storeCallbackPtr,
93+
sizeof(h_storeCallbackPtr)));
94+
95+
// Now associate the callbacks with the plan.
96+
cufftResult status = cufftXtSetCallback(fftPlan,
97+
(void **)&h_loadCallbackPtr,
98+
CUFFT_CB_LD_REAL,
99+
0);
100+
if (status == CUFFT_LICENSE_ERROR) {
101+
printf("This sample requires a valid license file.\n");
102+
printf("The file was either not found, out of date, or otherwise invalid.\n");
103+
exit(EXIT_FAILURE);
104+
} else {
105+
checkCudaErrors(status);
106+
}
107+
108+
checkCudaErrors(cufftXtSetCallback(fftPlan,
109+
(void **)&h_storeCallbackPtr,
110+
CUFFT_CB_ST_COMPLEX,
111+
(void **)&filter));
112+
113+
//create timers
114+
cudaEvent_t start, end;
115+
cudaEventCreate(&start);
116+
cudaEventCreate(&end);
117+
float elapsedTime;
118+
119+
printf("Running %d iterations\n", ITERATIONS);
120+
checkCudaErrors(cudaEventRecord(start, 0));
121+
122+
/*
123+
* The actual Computation
124+
*/
125+
126+
for(int i = 0; i < ITERATIONS; i++) {
127+
checkCudaErrors(cufftExecR2C(fftPlan, (cufftReal*)_8bit_signal, result));
128+
}
129+
130+
checkCudaErrors(cudaEventRecord(end, 0));
131+
checkCudaErrors(cudaEventSynchronize(end));
132+
checkCudaErrors(cudaEventElapsedTime(&elapsedTime, start, end));
133+
printf("Time for the FFT: %fms\n", elapsedTime);
134+
135+
//Verify correct result
136+
if(postprocess(reference, result, COMPLEX_SIGNAL_SIZE * BATCH_SIZE)) {
137+
printf("Verification successful.\n");
138+
} else {
139+
printf("!!! Verification Failed !!!\n");
140+
}
141+
142+
//Cleanup
143+
checkCudaErrors(cufftDestroy(fftPlan));
144+
145+
checkCudaErrors(cudaFree(_8bit_signal));
146+
checkCudaErrors(cudaFree(result));
147+
checkCudaErrors(cudaFree(filter));
148+
checkCudaErrors(cudaFree(reference));
149+
150+
//clean up driver state
151+
cudaDeviceReset();
152+
153+
printf("Done\n");
154+
155+
return 0;
156+
}

0 commit comments

Comments
 (0)