diff --git "a/CUDA\350\252\262\351\241\214/streamdivergencefft.cu" "b/CUDA\350\252\262\351\241\214/streamdivergencefft.cu" new file mode 100644 index 0000000000000000000000000000000000000000..749d52af2acc78004f1accca769fffdbf1f18c73 --- /dev/null +++ "b/CUDA\350\252\262\351\241\214/streamdivergencefft.cu" @@ -0,0 +1,74 @@ +#include +#include +#include +#define NX 512 +#define NY 512 +#define BLOCK_SIZE 4 + +// 複素数の構造体定義 +struct Complex { + float real; + float imag; +}; + +__global__ void InputHost(Complex *input) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + int j = blockIdx.y * blockDim.y + threadIdx.y; + if (i < NX && j < NY) { + input[i * NY + j].real = sinf(i + j); + input[i * NY + j].imag = cosf(i + j); + } +} + +int main() { + const int size = NX * NY; + const int complexSize = size * sizeof(Complex); + + cufftHandle plan; + + // ホスト上の出力データ + Complex hostOutput[size]; + + // デバイス上の入力データと出力データ + Complex *deviceInput; + Complex *deviceOutput; + + // デバイスメモリを確保 + cudaMalloc((void**)&deviceInput, complexSize); + cudaMalloc((void**)&deviceOutput, complexSize); + + // ストリームの生成 + cudaStream_t stream; + cudaStreamCreate(&stream); + + // グリッドとブロックの次元を設定 + dim3 blockDim(BLOCK_SIZE, BLOCK_SIZE); + dim3 gridDim((NX + blockDim.x - 1) / blockDim.x, (NY + blockDim.y - 1) / blockDim.y); + + // カーネルを実行して入力データ初期化 + InputHost<<>>(deviceInput); + + // 2次元FFTのプランを作成 + cufftPlan2d(&plan, NX, NY, CUFFT_C2C); + + // FFTを実行 + cufftExecC2C(plan, (cufftComplex*)deviceInput, (cufftComplex*)deviceOutput, CUFFT_FORWARD); + + // 出力データ転送 + cudaMemcpyAsync(hostOutput, deviceOutput, complexSize, cudaMemcpyDeviceToHost, stream); + + // 結果の出力 + for (int i = 0; i < NX; i++) { + for (int j = 0; j < NY; j++) { + int idx = i * NY + j; + printf("Output: %f + %fI\n", hostOutput[idx].real, hostOutput[idx].imag); + } + } + + cudaStreamDestroy(stream); + cufftDestroy(plan); + cudaFree(deviceInput); + cudaFree(deviceOutput); + + return 0; +} \ No newline at end of file