CUDA in Maya plugin

最近一直在想，是否可以在 Maya plugin 裡面使用 CUDA 來作運算。便趁著週末參考 CUDA SDK cppIntegration 的範例，依樣畫葫蘆地在 Maya yTwistNode 的這個範例上作個簡單的小實驗，把 y-twist deform 的計算移至 CUDA kernel function 裡…

環境設定

在撰寫 CUDA 的程式碼之前，我們必須先在原來的 cpp project 中新增一條 build rule…

將 CUDA SDK 附的 Cuda.rules 加入列表並啟用…
(為了方便起見，我們可先新增一個環境變數 CUDA_SDK_PATH 使其指向 CUDA SDK 的安裝位置。 而 Cuda.rules 所在的位置則是在 $(CUDA_SDK_PATH)common 裡面)

啟用 Cuda.rules 之後， project properties 便會出現一個 CUDA 的分類，一些相關的參數都可以在這裡作設定。

程式重構

環境設定完之後，我們新增一個 .cu 的檔案，實作與 cpp code 的溝通介面以及處理 y-twist 的 kernel function

#include
#include 

__global__ void  gpu_twist(float4 pts[], int len,
                           float magnitude, float env)
{
	for(int k=0; k<len; k++)
	{
		float ff = magnitude * pts[k].y * env;
		if( ff != 0.0f )
		{
			double cct= cosf(ff);
			double cst= sinf(ff);
			double tt = pts[k].x * cct - pts[k].z * cst;

			pts[k].z = pts[k].x * cst + pts[k].z * cct;
			pts[k].x = tt;
		}
	}
}

extern "C"
void	gpu_calc(float ptData[][4], unsigned int len,
                        float magnitude, float env)
{
	const unsigned int memSize = len * sizeof(float) * 4;
	float4 *d_data;

	cutilSafeCall( cudaMalloc((void**) &d_data, memSize) );

	// Copy data to gpu memory
	cutilSafeCall( cudaMemcpy(d_data, ptData, memSize,
                                     cudaMemcpyHostToDevice) );

	// Invoke the kernel with one thread
	gpu_twist<<<1,1>>>(d_data, len, magnitude, env);

	// Copy result from gpu memory to host memory
	cutilSafeCall( cudaMemcpy(ptData, d_data, memSize,
                                    cudaMemcpyDeviceToHost) );

	cudaFree(d_data);
}

回到原來的 yTwistNode.cpp 補上溝通介面的宣告，並對 deform function 作一點小幅修改…

extern "C"
void	gpu_calc(float ptData[][4], unsigned int len, float magnitude, float env);

MStatus
yTwist::deform( MDataBlock& block,
		MItGeometry& iter,
		const MMatrix& m,
		unsigned int multiIndex)
{
	MStatus status = MS::kSuccess;

	// determine the angle of the yTwist
	//
	MDataHandle angleData = block.inputValue(angle,&status);
	McheckErr(status, "Error getting angle data handlen");
	double magnitude = angleData.asDouble();

	// determine the envelope (this is a global scale factor)
	//
	MDataHandle envData = block.inputValue(envelope,&status);
	McheckErr(status, "Error getting envelope data handlen");
	float env = envData.asFloat();

	// { Modified block
	MPointArray ptArray;
	status = iter.allPositions(ptArray);

	// Get raw data
	float (*rawData)[4] = new float[ptArray.length()][4];
	ptArray.get(rawData);

	// Invoke cuda
	gpu_calc(rawData, ptArray.length(), magnitude, env);

	// Retrieve the data
	for(int i=0; i<ptArray.length(); ++i)
	{
		ptArray.set(rawData[i], i);
	}
	delete [] rawData;

	iter.setAllPositions(ptArray);
	// } End of modified

	return status;
}

接著在 project properties 補上一些相關設定:

• Additional Include Directories
  $(CUDA_INC_PATH); $(CUDA_SDK_PATH)commoninc
• Additional Library Directories
  $(CUDA_LIB_PATH); $(CUDA_SDK_PATH)commonlib
• Additional Dependencies
  cudart.lib cutil32D.lib
• CUDA 的 Additional Include Directories
  “$(CUDA_SDK_PATH)commoninc“
  Ps. 這裡的 path 記得要 quote 起來，不然會有 nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified 的錯誤訊息!!

最後的步驟當然是 build & execute 囉!! 至於執行效能，還真的如預期中的慢!? 這是因為 y-twist 其實也沒做多少事，花在 memcpy 的時間搞不好還比計算的時間多。加上我全部的資料都宣告在 global memory ， thread 也只用一個而已。所以執行效能較慢也是預料中的事情，不過經過這一次的實驗，總算大致熟悉了從 cpp project 引入 CUDA 的流程，也終於在 CUDA 的世界中心呼喊 Hello World 了!! 

Advertisement