Radeon(OpenCL)とVC2008ExpressでHelloWorld

GPU

前々からGPUプログラミングに興味がありましたが、環境が整わずなかなかチャレンジできない状況でした。


さて昨日、こんな記事を書いた後に思いついたのが、「そういえば少し前に格安Radeonを入手していた」という事実。


Intel HD GraphicsでGPUコンピューティング「OpenCLで暗号化復号化演算」 - .h2oのお気楽日記


さっそくAMD Accelerated Parallel Processing (APP) SDK(旧ATI Stream SDK)を使いOpenCLを試してみる事にした(できるだけコストをかけず無料の方向で)。

AMD APP technology is a set of advanced hardware and software technologies that enable AMD graphics processing cores (GPU), working in concert with the system’s x86 cores (CPU), to accelerate many applications beyond just graphics. This enables better balanced platforms capable of running demanding computing tasks faster than ever, and sets software developers on the path to optimize for AMD Accelerated Processing Units (APUs).

http://developer.amd.com/sdks/AMDAPPSDK/Pages/default.aspx
  • 環境
    • OS:WinXP
    • GPU:ATI Radeon5450
    • 開発環境:VC2008Express
      • VCの設定
        • インクルードファイル：C:\Program Files\AMD APP\include
        • ライブラリファイル：C:\Program Files\AMD APP\lib\x86
        • ライブラリ:OpenCL.libを追加

まずは、上記サイトでAPP SDKをダウンロードしセットアップします。


次に、VC上で"OpenCLSample"ソリューションを作成し、以下のコードを配置します。


ソースコードは大きく2つで構成されています"OpenCLSample.cpp"と"lesson1_kernels.cl"(VC上のプロジェクトに追加する)。

// OpenCLSample.cpp : コンソール アプリケーションのエントリ ポイントを定義します。
//
#include "stdafx.h"
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <iostream>
#include <string>
#include <iterator>

#include <utility>
#define __NO_STD_VECTOR // Use cl::vector and cl::string and
#define __NO_STD_STRING // not STL versions, more on this later
#include <CL/cl.hpp>

const std::string hw("Hello World\n");

inline void
checkErr(cl_int err, const char * name)
{
	if (err != CL_SUCCESS)
	{
		std::cerr << "ERROR: " << name
			<< " (" << err << ")" << std::endl;
		exit(EXIT_FAILURE);
	}
}

int main(void)
{
	cl_int err;

	cl::vector <cl::Platform> platformList;
	cl::Platform::get(&platformList);
	checkErr(
		platformList.size() != 0 ? CL_SUCCESS : -1, "cl::Platform::get");
	std::cerr << "Platform number is: " << platformList.size() << std::endl;

	cl::STRING_CLASS platformVendor;
	platformList[0].getInfo(CL_PLATFORM_VENDOR, &platformVendor);
	std::cerr << "Platform is by: " << platformVendor.c_str() << "\n";

	cl::STRING_CLASS platformVersion;
	platformList[0].getInfo(CL_PLATFORM_VERSION, &platformVersion);
	std::cerr << "Platform Version is by: " << platformVersion.c_str() << "\n";

	cl_context_properties cprops[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)(platformList[0])(), 0 };
	cl::Context context(
		CL_DEVICE_TYPE_GPU,
		cprops,
		NULL,
		NULL,
		&err);
	checkErr(err, "Context::Context()");

	char * outH = new char[hw.length() + 1];
	cl::Buffer outCL(
		context,
		CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
		hw.length() + 1,
		outH,
		&err);
	checkErr(err, "Buffer::Buffer()");

	cl::vector<cl::Device> devices;
	devices = context.getInfo<CL_CONTEXT_DEVICES>();
	checkErr(
		devices.size() > 0 ? CL_SUCCESS : -1, "devices.size() > 0");

	std::ifstream file("lesson1_kernels.cl");
	checkErr(file.is_open() ? CL_SUCCESS : -1, "lesson1_kernel.cl");

	std::string prog(
		std::istreambuf_iterator<char>(file),
		(std::istreambuf_iterator<char>()));

	cl::Program::Sources source(
		1,
		std::make_pair(prog.c_str(), prog.length() + 1));
	cl::Program program(context, source);
	err = program.build(devices, "");
	checkErr(
		file.is_open() ? CL_SUCCESS : -1, "Program::build()");

	cl::Kernel kernel(program, "hello", &err);
	checkErr(err, "Kernel::Kernel()");

	err = kernel.setArg(0, outCL);
	checkErr(err, "Kernel::setArg()");

	cl::CommandQueue queue(context, devices[0], 0, &err);
	checkErr(err, "CommandQueue::CommandQueue()");

	cl::Event event;
	err = queue.enqueueNDRangeKernel(
		kernel,
		cl::NullRange,
		cl::NDRange(hw.length() + 1),
		cl::NDRange(1, 1),
		NULL,
		&event);
	checkErr(err, "CommandQueue::enqueueNDRangeKernel()");

	event.wait();
	err = queue.enqueueReadBuffer(
		outCL,
		CL_TRUE,
		0,
		hw.length() + 1,
		outH);
	checkErr(err, "CommandQueue::enqueueReadBuffer()");
	std::cout << outH;

	Sleep(5000);

	return EXIT_SUCCESS;
}
// lesson1_kernels.cl
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
__constant char hw[] = "Hello World\n";
__kernel void hello(__global char * out)
{
	size_t tid = get_global_id(0);
	out[tid] = hw[tid];
}


あとは、ビルドし実行するだけ、ぽちっとな。

Platform number is: 1
Platform is by: Advanced Micro Devices, Inc.
Platform Version is by: OpenCL 1.1 AMD-APP-SDK-v2.4 (595.10)
Hello World
続行するには何かキーを押してください . . .

おお、演算はしていないけどなんか動いた感じ。


ふと思ったのは、AndroidやiPhone等スマートフォンでもOpenCLが求められる時代が来るかもしれないという事。


すでにGPU積んでるんだからフツーに使われる日も遠くないのかも…。

  • 今後やりたいこと

APP SDKのリファレンスではIntel Compiler/GCC系「MinGW(32bit or 64bit)」に対応済みという事で、MinGWを使いUNIX環境と親和性の高いコードを書きたいです。


今回は、こちらのページ（上記のコードや環境設定）を大変参考にさせて戴きました。
めも ATI Stream SDKでopenCLその４